[all-commits] [llvm/llvm-project] 125143: Initial commit of MIOpen dialect.
Wen-Heng (Jack) Chung via All-commits
all-commits at lists.llvm.org
Thu Oct 22 13:20:19 PDT 2020
Branch: refs/heads/miopen-dialect-tweak-perf-jenkins
Home: https://github.com/llvm/llvm-project
Commit: 125143fddbefacd6f10e70fcd2f1db52ddaaeb16
https://github.com/llvm/llvm-project/commit/125143fddbefacd6f10e70fcd2f1db52ddaaeb16
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/CMakeLists.txt
A mlir/include/mlir/Dialect/MIOpenOps/CMakeLists.txt
A mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h
M mlir/lib/Dialect/CMakeLists.txt
A mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
A mlir/lib/Dialect/MIOpenOps/DialectRegistration.cpp
A mlir/lib/Dialect/MIOpenOps/MIOpenOps.cpp
Log Message:
-----------
Initial commit of MIOpen dialect.
Commit: 4600d68a403647ea41eca5c62616d7abd77d9281
https://github.com/llvm/llvm-project/commit/4600d68a403647ea41eca5c62616d7abd77d9281
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h
A mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
M mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
M mlir/lib/Dialect/MIOpenOps/MIOpenOps.cpp
A mlir/test/Dialect/MIOpen/ops.mlir
Log Message:
-----------
Initial commit to add ops into MIOpen dialect.
- conv2d
- transform
- gridwise_gemm
add dummy parse / print / verify logic.
add dummy test.
Commit: ff39c4c709ac1603d21f7baab75dbfbb13ae6fbc
https://github.com/llvm/llvm-project/commit/ff39c4c709ac1603d21f7baab75dbfbb13ae6fbc
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
M mlir/lib/Dialect/MIOpenOps/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops.mlir
Log Message:
-----------
Add parse / print logic to MIOpen ops.
Revise test cases along the way.
Commit: 9bc39d9e6f9e7046f99dc105c093c78fc8031011
https://github.com/llvm/llvm-project/commit/9bc39d9e6f9e7046f99dc105c093c78fc8031011
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
A mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
M mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
A mlir/lib/Dialect/MIOpenOps/CppOutput/CMakeLists.txt
A mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
A mlir/test/Dialect/MIOpen/CppOutput/miopencpp.mlir
Log Message:
-----------
Initial commit to introduce MLIR MIOpen dialect -> MIOpen C++ translation.
Commit: 05d19a7eb4ea27d4d0e7989b145c2f7458fff54a
https://github.com/llvm/llvm-project/commit/05d19a7eb4ea27d4d0e7989b145c2f7458fff54a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
M mlir/lib/Dialect/MIOpenOps/CppOutput/CMakeLists.txt
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
R mlir/test/Dialect/MIOpen/CppOutput/miopencpp.mlir
A mlir/test/Dialect/MIOpen/CppOutput/transformed.mlir
Log Message:
-----------
Add Op traversing logic into MIOpen dialect -> C++ translator.
Commit: f0500d18ee6f0072487526e53fc9a7bd1c1d7235
https://github.com/llvm/llvm-project/commit/f0500d18ee6f0072487526e53fc9a7bd1c1d7235
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
R mlir/test/Dialect/MIOpen/CppOutput/transformed.mlir
A mlir/test/Dialect/MIOpen/translate.mlir
Log Message:
-----------
Add Op traversing logic into MIOpen dialect -> C++ header translator.
Commit: 40ca3b593063c2628082dbdff6b31b8e728cb926
https://github.com/llvm/llvm-project/commit/40ca3b593063c2628082dbdff6b31b8e728cb926
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/test/Dialect/MIOpen/ops.mlir
Log Message:
-----------
Amend op parsing test case.
Added the following attributes:
- source_layout
- output_layout
- gridwise_gemm_argument_position
Commit: 7054cfc71ac450ab5ac9ee505d1096bb1252f9c6
https://github.com/llvm/llvm-project/commit/7054cfc71ac450ab5ac9ee505d1096bb1252f9c6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
Log Message:
-----------
Fix translation logic.
Commit: e587a8a23c9812610dd1e79a3e1211e1f4d8aba5
https://github.com/llvm/llvm-project/commit/e587a8a23c9812610dd1e79a3e1211e1f4d8aba5
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
A mlir/include/mlir/Dialect/MIOpenOps/Passes.h
A mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
A mlir/test/Dialect/MIOpen/lowering.mlir
Log Message:
-----------
Inital commit to add MIOpen Conv2D to Transform and GridwiseGemm transform pass.
Commit: 1c3be7ec0838048d0a5a8f2ebf3dfa5e831370cd
https://github.com/llvm/llvm-project/commit/1c3be7ec0838048d0a5a8f2ebf3dfa5e831370cd
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
M mlir/test/Dialect/MIOpen/lowering.mlir
Log Message:
-----------
Add Op transform logic. Improve Op translate logic.
Revise tests.
Commit: 0ac976c8c13175d7f744156dba3302f7d1457073
https://github.com/llvm/llvm-project/commit/0ac976c8c13175d7f744156dba3302f7d1457073
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
M mlir/test/Dialect/MIOpen/translate.mlir
Log Message:
-----------
Change emitted names to fixed "mlir".
Commit: fe71429c14c28ee85ff94a528c59b6a1436b33ac
https://github.com/llvm/llvm-project/commit/fe71429c14c28ee85ff94a528c59b6a1436b33ac
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
Log Message:
-----------
Fix input / weight tensor translation logic in Cpp preamble.
Commit: 01659f382d0ab554d59bfd793a0f9153613c7e20
https://github.com/llvm/llvm-project/commit/01659f382d0ab554d59bfd793a0f9153613c7e20
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
A mlir/test/Dialect/MIOpen/lowering_kcyx_nchw_nkhw.mlir
A mlir/test/Dialect/MIOpen/lowering_kyxc_nhwc_nhwk.mlir
Log Message:
-----------
Generalized op transformation logic for weight tensor.
Add test cases.
Commit: a1e3fec79420164b7cd398872d525f03c4436e96
https://github.com/llvm/llvm-project/commit/a1e3fec79420164b7cd398872d525f03c4436e96
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
A mlir/test/Dialect/MIOpen/lowering_ckyx_cnhw_knhw.mlir
A mlir/test/Dialect/MIOpen/lowering_cyxk_chwn_khwn.mlir
A mlir/test/Dialect/MIOpen/lowering_cyxk_cnhw_knhw.mlir
Log Message:
-----------
Generalized op transformation logic for output tensor.
Add more op lowering test cases.
Commit: 346c9136bf1fcf9336b971b52059f6698145d9cd
https://github.com/llvm/llvm-project/commit/346c9136bf1fcf9336b971b52059f6698145d9cd
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
Log Message:
-----------
Fix stride translation logic.
Commit: e404d0d39c71d8019abd59d43fc31691387a6233
https://github.com/llvm/llvm-project/commit/e404d0d39c71d8019abd59d43fc31691387a6233
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Fix comment.
Commit: 545a31123df8fc39dbdb06b634ba36d450553bc1
https://github.com/llvm/llvm-project/commit/545a31123df8fc39dbdb06b634ba36d450553bc1
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Generalized op transformation logic for input tensor step 1: padded input.
Commit: 0214f2772ebf13893306df536f1a3c39f13f0500
https://github.com/llvm/llvm-project/commit/0214f2772ebf13893306df536f1a3c39f13f0500
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Generalized op transformation logic for input tensor step 2: embedded input.
Commit: 6f6195448feee799533c878d0ad898a7a6f25c9d
https://github.com/llvm/llvm-project/commit/6f6195448feee799533c878d0ad898a7a6f25c9d
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Generalized op transformation logic for input tensor step 3: transformed input.
Via dimension merging.
Commit: 65d9e26c995ad05be041d036f301a669466158f8
https://github.com/llvm/llvm-project/commit/65d9e26c995ad05be041d036f301a669466158f8
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
Log Message:
-----------
Fix parameters in translator logic.
Commit: b2ba53adda747b0c9d17a520ce1b20e6aa81bf02
https://github.com/llvm/llvm-project/commit/b2ba53adda747b0c9d17a520ce1b20e6aa81bf02
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Code refactor. Extract attribute consturction routines.
Commit: f5e93707ee93dbc95cd11d7e664b57244478d519
https://github.com/llvm/llvm-project/commit/f5e93707ee93dbc95cd11d7e664b57244478d519
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Further code refactor. Move attribute construction logic to tablegen.
Commit: de82a3f711e747c9d381e3166ef654a7c5346b1e
https://github.com/llvm/llvm-project/commit/de82a3f711e747c9d381e3166ef654a7c5346b1e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Further code refactor for ArrayAttr and DictionaryAttr.
Commit: 2070b3470eed31bd78bbd33f815596f10e36c6e5
https://github.com/llvm/llvm-project/commit/2070b3470eed31bd78bbd33f815596f10e36c6e5
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Further refactor. Use Builder API.
Commit: 02d8abac8423dd00ce379036966018869fbc8a24
https://github.com/llvm/llvm-project/commit/02d8abac8423dd00ce379036966018869fbc8a24
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Further refactor.
Commit: c9808d0c81139ecbd239a8aa00840c60232f305f
https://github.com/llvm/llvm-project/commit/c9808d0c81139ecbd239a8aa00840c60232f305f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
A mlir/lib/Dialect/MIOpenOps/Driver/CMakeLists.txt
A mlir/lib/Dialect/MIOpenOps/Driver/mlir-miopen-driver.cpp
Log Message:
-----------
Introduce a new target for high-level op generation: mlir-miopen-driver.
Commit: 804e1e275e929467d48d456f9e55e32c6387fbe7
https://github.com/llvm/llvm-project/commit/804e1e275e929467d48d456f9e55e32c6387fbe7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/Driver/mlir-miopen-driver.cpp
Log Message:
-----------
Implement main logic.
Commit: 3ac107afc07d8cb279692f3a8d87c0d4863203fb
https://github.com/llvm/llvm-project/commit/3ac107afc07d8cb279692f3a8d87c0d4863203fb
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/Driver/mlir-miopen-driver.cpp
Log Message:
-----------
Use proper file output logic.
Commit: 7cc5ed8fdd7b69ac85857509d17a5b88a43ea80f
https://github.com/llvm/llvm-project/commit/7cc5ed8fdd7b69ac85857509d17a5b88a43ea80f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
M mlir/lib/Dialect/MIOpenOps/CppOutput/CMakeLists.txt
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
A mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
A mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Initial commit to support multiple tranlation targets.
Commit: 02578ac2b29690430620ac630fcfde48273e9fda
https://github.com/llvm/llvm-project/commit/02578ac2b29690430620ac630fcfde48273e9fda
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Add Op traversing logic into MIOpen dialect -> C++ translator with XDLOPS support.
Commit: 7cb2738a2e7bdfb7f33b20cb640543c053438bdd
https://github.com/llvm/llvm-project/commit/7cb2738a2e7bdfb7f33b20cb640543c053438bdd
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Add Op traversing logic into MIOpen dialect -> C++ header translator with XDLOPS support.
Commit: aa8c5ed730e1db2df717b22bff253c4fa4a0e09b
https://github.com/llvm/llvm-project/commit/aa8c5ed730e1db2df717b22bff253c4fa4a0e09b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/Driver/mlir-miopen-driver.cpp
Log Message:
-----------
Revise mlir-miopen-driver command line arguments.
Commit: 70164ac1da4f8159dec07860bdb31ed5a83276be
https://github.com/llvm/llvm-project/commit/70164ac1da4f8159dec07860bdb31ed5a83276be
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/Driver/mlir-miopen-driver.cpp
Log Message:
-----------
Add new --pop flag for mlir-miopen-driver to populate default values.
Commit: 6fc00bb476049b654142fb5120f1319ecce8acdd
https://github.com/llvm/llvm-project/commit/6fc00bb476049b654142fb5120f1319ecce8acdd
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops.mlir
Log Message:
-----------
Add logic to set attributes for miopen.gridwise_gemm op.
Commit: 54de6dee0b751ab5e5cf16bfb0c43f9c5bdc5a74
https://github.com/llvm/llvm-project/commit/54de6dee0b751ab5e5cf16bfb0c43f9c5bdc5a74
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/Driver/mlir-miopen-driver.cpp
Log Message:
-----------
Revise mlir-miopen-driver command line arguments.
Commit: 8ed1b97c2362776c8a4781038291a4252acaf30b
https://github.com/llvm/llvm-project/commit/8ed1b97c2362776c8a4781038291a4252acaf30b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Initial commit support producing C++ compilation flags.
Commit: 93b50ee66ac753547e517fbfb7be69a54109f086
https://github.com/llvm/llvm-project/commit/93b50ee66ac753547e517fbfb7be69a54109f086
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Implement logic to emit compilation flags for non-XDLOPS convolution.
Commit: a1177802afd431a3a040cbc60fe71f6fd06c1acb
https://github.com/llvm/llvm-project/commit/a1177802afd431a3a040cbc60fe71f6fd06c1acb
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Implement logic to emit compilation flags for non-XDLOPS convolution.
Commit: 3ab712af61b0689e917f3de5422374fbb34679b9
https://github.com/llvm/llvm-project/commit/3ab712af61b0689e917f3de5422374fbb34679b9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Change the order of emitting compilation flags.
Distinguish those from convolution configs and other heuristic-based or
hard-coded ones.
Commit: 5419bca9f1f7c25752a99ca60ab682fb80febe90
https://github.com/llvm/llvm-project/commit/5419bca9f1f7c25752a99ca60ab682fb80febe90
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Extract tunable parameters into a separate structure.
Commit: bb64bbc0e5b7104c51806b9685acfbcabcd02ee1
https://github.com/llvm/llvm-project/commit/bb64bbc0e5b7104c51806b9685acfbcabcd02ee1
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Start adding YAML I/O routines for tunable parameters.
Commit: c07ebe849e4f6d60b3571f2d31b71afc657b907f
https://github.com/llvm/llvm-project/commit/c07ebe849e4f6d60b3571f2d31b71afc657b907f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Extract YAML I/O routines to a common header.
Commit: 1b7759afa07e89c47721620478663e09277b15c9
https://github.com/llvm/llvm-project/commit/1b7759afa07e89c47721620478663e09277b15c9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Emit values derived from tunable parameters.
Commit: 3eeedc939d2f4e937f5506f22cb5a0d301f4359b
https://github.com/llvm/llvm-project/commit/3eeedc939d2f4e937f5506f22cb5a0d301f4359b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
M mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Be able to load from YAML and populate YAML configs.
Commit: bd32b79c491fe95f4ce9b62db5c40ec12108d837
https://github.com/llvm/llvm-project/commit/bd32b79c491fe95f4ce9b62db5c40ec12108d837
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
M mlir/test/Dialect/MIOpen/ops.mlir
Log Message:
-----------
Various minor fixes.
Commit: 9dadd6e82ecdcda516c69b7f8df0ace95074a835
https://github.com/llvm/llvm-project/commit/9dadd6e82ecdcda516c69b7f8df0ace95074a835
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Reorder tunable parameters a bit.
Next commit would add logic to determine vectorization dimension and vectorize width.
Commit: 031f4fdd8f6a3c4bde7c7d57081a3147217c40e0
https://github.com/llvm/llvm-project/commit/031f4fdd8f6a3c4bde7c7d57081a3147217c40e0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Fix stack corruption issue.
Commit: d888df0c272ec447cf5600218691e4648d6955a0
https://github.com/llvm/llvm-project/commit/d888df0c272ec447cf5600218691e4648d6955a0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Initial logic to determine vectorization dimension and length.
Commit: 0df739ab5dc436682365c5fed985f4052551dc94
https://github.com/llvm/llvm-project/commit/0df739ab5dc436682365c5fed985f4052551dc94
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Add logic to fix header emission logic on filter tensor vectorizable dimension.
Commit: a09cbfd9070462266b63ad929aee9ec188e64b26
https://github.com/llvm/llvm-project/commit/a09cbfd9070462266b63ad929aee9ec188e64b26
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Add logic to fix header emission logic on input tensor vectorizable dimension.
Commit: c00e2b70d8d1dd15ab2597a14284d5237739fba1
https://github.com/llvm/llvm-project/commit/c00e2b70d8d1dd15ab2597a14284d5237739fba1
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/test/Dialect/MIOpen/translate.mlir
Log Message:
-----------
Fix unit test.
Commit: f0c88afbf5b39f7f1c0d19180e4d444c1ca16352
https://github.com/llvm/llvm-project/commit/f0c88afbf5b39f7f1c0d19180e4d444c1ca16352
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Fine tune default parameters.
Commit: 266a9588f638be695c3ff13a6732658deaa1c612
https://github.com/llvm/llvm-project/commit/266a9588f638be695c3ff13a6732658deaa1c612
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Improve vectorization computation logic.
Commit: 97d2b37bfb38e71c9c8abdfd9144599b715f9d02
https://github.com/llvm/llvm-project/commit/97d2b37bfb38e71c9c8abdfd9144599b715f9d02
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Minor refactoring and comment polish.
Commit: 3a68376ef3abcde794b686851248e203a6ade152
https://github.com/llvm/llvm-project/commit/3a68376ef3abcde794b686851248e203a6ade152
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Add logic to consider right padding.
Commit: 716cd176753816bd0cc4b9c377d540e85bcadd2d
https://github.com/llvm/llvm-project/commit/716cd176753816bd0cc4b9c377d540e85bcadd2d
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Revise XDLOPS C++ kernel emitter logic to reflect latest changes in MIOpen.
Commit: f35b95e00a35dcc9b1eeff03648b3135b8b93e29
https://github.com/llvm/llvm-project/commit/f35b95e00a35dcc9b1eeff03648b3135b8b93e29
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Revise non-XDLOPS C++ kernel emitter logic to reflect latest changes in MIOpen.
Commit: a39f68b9b1f94a1f2b398d66497d939229cd95c0
https://github.com/llvm/llvm-project/commit/a39f68b9b1f94a1f2b398d66497d939229cd95c0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Revise XDLOPS C++ kernel emitter logic to reflect latest changes in MIOpen.
Commit: 5ffec51b4c21a9ba56d9969d1b081a5a46245192
https://github.com/llvm/llvm-project/commit/5ffec51b4c21a9ba56d9969d1b081a5a46245192
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Revise output tensor transformation logic.
Commit: fba0991448fbc9ff16a5c68bf9db1d63c89b1dcf
https://github.com/llvm/llvm-project/commit/fba0991448fbc9ff16a5c68bf9db1d63c89b1dcf
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Heavily revise XDLOPS C++ code emission logic.
Commit: 6875cc07924df96d9ba385c87f082e2d419cf408
https://github.com/llvm/llvm-project/commit/6875cc07924df96d9ba385c87f082e2d419cf408
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
M mlir/lib/Dialect/MIOpenOps/Driver/mlir-miopen-driver.cpp
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
M mlir/lib/Dialect/MIOpenOps/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops.mlir
Log Message:
-----------
Adding conv2d_backward_data to parser
* Updated table-gen to generate the new conv2d_backward_data high level
operation
* Updated mlir-miopen-driver to branch based on conv_flavor
* Added kernel_algorithm field to existing lowering step
* Added conv2d_backward_data test case
Commit: c64f0addd96014649b7030f29828348ba0844259
https://github.com/llvm/llvm-project/commit/c64f0addd96014649b7030f29828348ba0844259
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/Driver/mlir-miopen-driver.cpp
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Address review feedbacks
Commit: 7eed23fd099c5dc93bedd3ec442d9339883ec438
https://github.com/llvm/llvm-project/commit/7eed23fd099c5dc93bedd3ec442d9339883ec438
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Revise XDLOPS vectorization flag emission logic.
- Port logic from MIOpen.
- Consider generic layouts.
Commit: 587540d675d75320797fa1e7337d03b47ff08950
https://github.com/llvm/llvm-project/commit/587540d675d75320797fa1e7337d03b47ff08950
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
XXX Optimization: use Unfold to replace Merge on XDLOPS for Weight tensor.
Commit: ff3096e04996ef07f08530fa400cc6f77929abdd
https://github.com/llvm/llvm-project/commit/ff3096e04996ef07f08530fa400cc6f77929abdd
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Additional vectorization logic for Input tensor (Matrix B).
Commit: 3139c58e7b673003b2ef2fb3371a8872454fe34c
https://github.com/llvm/llvm-project/commit/3139c58e7b673003b2ef2fb3371a8872454fe34c
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
A mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Adding lowering for conv_bwd_data
* templatize the conv and conv_bwd_data common logic into header file
* Use argument_fields to capture differences between conv operations
Commit: 6ac91474030b7e86063bf689e0a75d0bbb0264f4
https://github.com/llvm/llvm-project/commit/6ac91474030b7e86063bf689e0a75d0bbb0264f4
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
Log Message:
-----------
Reapply unfold optimization in 05c6f39367e09a7cbce14c9f55005bc83aad79ba
Commit: f4783f633b2e7658a3b3149b69b2c18ac1a54773
https://github.com/llvm/llvm-project/commit/f4783f633b2e7658a3b3149b69b2c18ac1a54773
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/test/Dialect/MIOpen/lowering.mlir
M mlir/test/Dialect/MIOpen/lowering_ckyx_cnhw_knhw.mlir
M mlir/test/Dialect/MIOpen/lowering_cyxk_chwn_khwn.mlir
M mlir/test/Dialect/MIOpen/lowering_cyxk_cnhw_knhw.mlir
M mlir/test/Dialect/MIOpen/lowering_kcyx_nchw_nkhw.mlir
M mlir/test/Dialect/MIOpen/lowering_kyxc_nhwc_nhwk.mlir
Log Message:
-----------
Addressing review feedback and adding test cases
Commit: b2f54d0b20f341094d7a253e0881987771bed9f5
https://github.com/llvm/llvm-project/commit/b2f54d0b20f341094d7a253e0881987771bed9f5
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
R mlir/test/Dialect/MIOpen/lowering.mlir
R mlir/test/Dialect/MIOpen/lowering_ckyx_cnhw_knhw.mlir
R mlir/test/Dialect/MIOpen/lowering_cyxk_chwn_khwn.mlir
R mlir/test/Dialect/MIOpen/lowering_cyxk_cnhw_knhw.mlir
A mlir/test/Dialect/MIOpen/lowering_filter_tensor_ckyx_cnhw_knhw.mlir
A mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_position_cyxk_chwn_khwn.mlir
A mlir/test/Dialect/MIOpen/lowering_input_tensor_cyxk_cnhw_hnhw.mlir
R mlir/test/Dialect/MIOpen/lowering_kcyx_nchw_nkhw.mlir
R mlir/test/Dialect/MIOpen/lowering_kyxc_nhwc_nhwk.mlir
A mlir/test/Dialect/MIOpen/lowering_memref_kcyx_nchw_nkhw.mlir
A mlir/test/Dialect/MIOpen/lowering_output_tensor_kyxc_nhwc_nhwk.mlir
A mlir/test/Dialect/MIOpen/lowering_top_level.mlir
Log Message:
-----------
Rename unit test files to include test purpose
Commit: c2aea2eda8ed77c96198b6c7e53cdc3fa53b7062
https://github.com/llvm/llvm-project/commit/c2aea2eda8ed77c96198b6c7e53cdc3fa53b7062
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h
Log Message:
-----------
Moving ConvOpType to MIOpenOps.h to make it shared
Commit: 79587000d00472073341cf36bc62db091ce36f8e
https://github.com/llvm/llvm-project/commit/79587000d00472073341cf36bc62db091ce36f8e
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Removing unused layoutstring from arguments
Commit: 5d809a508bbb7d8bff2e6b1028b2d0608cf04d37
https://github.com/llvm/llvm-project/commit/5d809a508bbb7d8bff2e6b1028b2d0608cf04d37
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
R mlir/test/Dialect/MIOpen/translate.mlir
A mlir/test/Dialect/MIOpen/translate_bwd_data_kcyx_nchw_nkhw.mlir
A mlir/test/Dialect/MIOpen/translate_forward_kcyx_nchw_nkhw.mlir
Log Message:
-----------
Adding conv2d_bwd_data cpp translation
* Added kernel_algorithm field to lowering
* Added kernel_algorithm check to determine ConvOpType
* Updated constexpr string to be format string in different ConvOpType
* Added conv2d_bwd_data translate cpp test
Commit: 8b0a812af5f83635b3036cf315d96874ec3ff170
https://github.com/llvm/llvm-project/commit/8b0a812af5f83635b3036cf315d96874ec3ff170
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
M mlir/lib/Dialect/MIOpenOps/MIOpenOps.cpp
A mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Add new operations to help realize lowering gridwise GEMM.
Commit: c1dfd987f2aff11346fe40340bd2431958d2fa6e
https://github.com/llvm/llvm-project/commit/c1dfd987f2aff11346fe40340bd2431958d2fa6e
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/test/Dialect/MIOpen/translate_bwd_data_kcyx_nchw_nkhw.mlir
Log Message:
-----------
Adding Conv2d_bwd_data header translation
* Updated logic to determine vectorization
* Branched off cpp translation according to opType
* Added hpp test to translation
Commit: eb2c58376bdf67db77b77c5409fa746ed12fc9bb
https://github.com/llvm/llvm-project/commit/eb2c58376bdf67db77b77c5409fa746ed12fc9bb
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
A mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_level.mlir
Log Message:
-----------
Initial commit of skeleton for gridwise_gemm_ex lowering logic.
Commit: 8669178dc4216a9c1a21c15cea952822113e9b98
https://github.com/llvm/llvm-project/commit/8669178dc4216a9c1a21c15cea952822113e9b98
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
Log Message:
-----------
WIP lowering logic for gridwise GEMM op.
Commit: be5e21481bd93a8f4cd282c07d7af51cc4f6e0d9
https://github.com/llvm/llvm-project/commit/be5e21481bd93a8f4cd282c07d7af51cc4f6e0d9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
WIP lowering logic for blockwise GEMM op.
Commit: d8eb253697914d7946e009745c3cd2fbd11ccd9e
https://github.com/llvm/llvm-project/commit/d8eb253697914d7946e009745c3cd2fbd11ccd9e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
M mlir/lib/Dialect/MIOpenOps/MIOpenOps.cpp
Log Message:
-----------
Add threadwise GEMM op and its lowering logic.
Commit: a1cf7f56f71f838d0f96d6ea82eec56c35a2fbb2
https://github.com/llvm/llvm-project/commit/a1cf7f56f71f838d0f96d6ea82eec56c35a2fbb2
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
M mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_level.mlir
Log Message:
-----------
Add lowering logic from BlockwiseCopyOp to ThreadwiseCopyOp.
Commit: c44baf73bf01ce7d0bf66d36a550b55b4c2fb8ea
https://github.com/llvm/llvm-project/commit/c44baf73bf01ce7d0bf66d36a550b55b4c2fb8ea
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
Log Message:
-----------
Add more hints on how attributes should be populated.
Commit: 2c639e06829ed0ba03f11772ff4c58f085e69491
https://github.com/llvm/llvm-project/commit/2c639e06829ed0ba03f11772ff4c58f085e69491
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/Passes.h
M mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
A mlir/lib/Dialect/MIOpenOps/Transform/AffineTransforms.cpp
A mlir/lib/Dialect/MIOpenOps/Transform/CMakeLists.txt
A mlir/test/Dialect/MIOpen/lowering_transform.mlir
A mlir/test/Dialect/MIOpen/lowering_transform2.mlir
M mlir/test/Dialect/MIOpen/ops.mlir
Log Message:
-----------
add Affine Transform Pass which uses AffineMap.
implemented:
- PassThrough
- Pad
- Merge
- Unfold
- Embed
Commit: a1baea9c10a7d52b5ad8849df3d6868aa08585e4
https://github.com/llvm/llvm-project/commit/a1baea9c10a7d52b5ad8849df3d6868aa08585e4
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
Log Message:
-----------
Improve shape computation logic in lowering logic.
implemented:
- filter tensor
- output tensor
- input tensor
TBD:
- padding parameters for "pad" operation in input tensor.
Commit: 54287abc12286be3676eb70944202b77af85f0fc
https://github.com/llvm/llvm-project/commit/54287abc12286be3676eb70944202b77af85f0fc
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/Passes.h
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Break operation lowering process into multiple steps.
- step 1 : conv -> transform gridwise operations
- (annotate affine-transformation pass here)
- step 2 : gridwise operations -> blockwise operations
- step 3: blockwise operations -> threadwise operations
- (TBD) step 4 : threadwise operations -> LLVM dialect
Commit: 47917d61293a5f8c6a464b091c530bfb819844cc
https://github.com/llvm/llvm-project/commit/47917d61293a5f8c6a464b091c530bfb819844cc
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
Log Message:
-----------
Hook GridwiseGemmRewrite up, prepare to ditch GridwiseGemmEx.
Commit: c388b596da4bb8d19412e5394d7e0b371392f880
https://github.com/llvm/llvm-project/commit/c388b596da4bb8d19412e5394d7e0b371392f880
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/Passes.h
M mlir/lib/Dialect/MIOpenOps/Transform/AffineTransforms.cpp
A mlir/lib/Dialect/MIOpenOps/Transform/AffixTuningParameters.cpp
M mlir/lib/Dialect/MIOpenOps/Transform/CMakeLists.txt
Log Message:
-----------
Populate some hard-coded tuning parameters to gridwise_gemm.
This allows computation in miopen-lowering-step2 and miopen-lowering-step3 possible.
Commit: a454f57ebeec532f1132a17c5f104b1e2b32534e
https://github.com/llvm/llvm-project/commit/a454f57ebeec532f1132a17c5f104b1e2b32534e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
Log Message:
-----------
Compute LDS sizes required for GridwiseGemmOp from tuning parameters.
Commit: 0ca2533c348fc33aab4e370f3a88dba384d9d8d9
https://github.com/llvm/llvm-project/commit/0ca2533c348fc33aab4e370f3a88dba384d9d8d9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpenOps/Transform/AffixTuningParameters.cpp
Log Message:
-----------
Compute register sizes for Matrix C on VGPRs.
Commit: 52a70a291eefda3fbfe67520653e1c0fb6850f86
https://github.com/llvm/llvm-project/commit/52a70a291eefda3fbfe67520653e1c0fb6850f86
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
Log Message:
-----------
Compute loop iterations in GridwiseGemmOp lowering.
Commit: 37658a47186bfa49e8ec1782c740a10ba594f2e0
https://github.com/llvm/llvm-project/commit/37658a47186bfa49e8ec1782c740a10ba594f2e0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
Log Message:
-----------
Compute BlockwiseGemmOp and BlockwiseCopyOp attributes from GridwiseGemmOp attributes.
Commit: 34fd4e8c5de1438d0473e8bb059c222689d44b5d
https://github.com/llvm/llvm-project/commit/34fd4e8c5de1438d0473e8bb059c222689d44b5d
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
A mlir/include/mlir/Dialect/MIOpenOps/ConvertMIOpenOpsToLLVM.h
M mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
A mlir/lib/Dialect/MIOpenOps/LLVMOutput/CMakeLists.txt
A mlir/lib/Dialect/MIOpenOps/LLVMOutput/ConvertMIOpenOpsToLLVM.cpp
Log Message:
-----------
Initial commit to introduce MIOpen dialect -> LLVM dialect passes.
Basically this pass currently does nothing instead of removing all MIOpen ops
from its enclosing function.
Commit: 5d0af0e802e4fd930b8a2b6630ffc666212e87a1
https://github.com/llvm/llvm-project/commit/5d0af0e802e4fd930b8a2b6630ffc666212e87a1
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LLVMOutput/ConvertMIOpenOpsToLLVM.cpp
Log Message:
-----------
Lower miopen.gpu_alloc op to LLVM.
- build alloca op for VGPRs.
- TBD: how to explicit make it address apce 5?
- build global op for LDS.
Commit: 10f4843585104ff75dd7b34454be1556164e981d
https://github.com/llvm/llvm-project/commit/10f4843585104ff75dd7b34454be1556164e981d
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LLVMOutput/ConvertMIOpenOpsToLLVM.cpp
Log Message:
-----------
Lower miopen.subview op to LLVM.
- Use LLVM GEP.
Commit: d87f8d1c52c39523157c5492b9e5c072b3987c39
https://github.com/llvm/llvm-project/commit/d87f8d1c52c39523157c5492b9e5c072b3987c39
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/ConvertMIOpenOpsToLLVM.h
M mlir/lib/Dialect/MIOpenOps/LLVMOutput/ConvertMIOpenOpsToLLVM.cpp
Log Message:
-----------
Lower miopen.lds_barrier to Std.
Commit: 5064128f35b1d6141df589a64d2ee2f37c556e4f
https://github.com/llvm/llvm-project/commit/5064128f35b1d6141df589a64d2ee2f37c556e4f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
R mlir/include/mlir/Dialect/MIOpenOps/ConvertMIOpenOpsToLLVM.h
A mlir/include/mlir/Dialect/MIOpenOps/ConvertMIOpenOpsToStd.h
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpenOps/LLVMOutput/CMakeLists.txt
R mlir/lib/Dialect/MIOpenOps/LLVMOutput/ConvertMIOpenOpsToLLVM.cpp
A mlir/lib/Dialect/MIOpenOps/LLVMOutput/ConvertMIOpenOpsToStd.cpp
M mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
Log Message:
-----------
Lower miopen.fill to Std. Also change miopen.gpu_alloc / miopen.subview to Std.
Commit: d1e7877d4b60920aecebc00637e99970f87d9ffb
https://github.com/llvm/llvm-project/commit/d1e7877d4b60920aecebc00637e99970f87d9ffb
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
M mlir/lib/Dialect/MIOpenOps/LLVMOutput/ConvertMIOpenOpsToStd.cpp
Log Message:
-----------
Make blockwise_gemm take 1D memref so we don't need std.view op.
Commit: 353a7871aa368ac7f66dec20f42eb86fafe6ed53
https://github.com/llvm/llvm-project/commit/353a7871aa368ac7f66dec20f42eb86fafe6ed53
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/LLVMOutput/ConvertMIOpenOpsToStd.cpp
Log Message:
-----------
Assign AffineExpr when lowering miopen.subview to std.subview.
Commit: f3d275e5412b33c8ceeef5d264622ddc6f51530d
https://github.com/llvm/llvm-project/commit/f3d275e5412b33c8ceeef5d264622ddc6f51530d
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
R mlir/lib/Dialect/MIOpenOps/LLVMOutput/CMakeLists.txt
R mlir/lib/Dialect/MIOpenOps/LLVMOutput/ConvertMIOpenOpsToStd.cpp
M mlir/lib/Dialect/MIOpenOps/Transform/CMakeLists.txt
A mlir/lib/Dialect/MIOpenOps/Transform/ConvertMIOpenOpsToStd.cpp
Log Message:
-----------
Remove LLVMOutput directory as we now lower to Std.
Commit: 1ba8687228040dd865a9da51929338479c10a27c
https://github.com/llvm/llvm-project/commit/1ba8687228040dd865a9da51929338479c10a27c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/Transform/AffineTransforms.cpp
Log Message:
-----------
Initial step to refactor -miopen-affine-transform pass.
There would be at most 3 affine maps be produced for an miopen.transform op.
- index
- index diff
- offset diff
Commit: 4ec5022456be5769fa413ca4b6c890928bff6eee
https://github.com/llvm/llvm-project/commit/4ec5022456be5769fa413ca4b6c890928bff6eee
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
A mlir/examples/prototypes/affine_test.mlir
A mlir/examples/prototypes/blah.mlir
A mlir/examples/prototypes/blah2.mlir
A mlir/examples/prototypes/blah3.mlir
A mlir/examples/prototypes/memref_addrspace_test.mlir
A mlir/examples/prototypes/prototype_threadwise_copy.mlir
A mlir/examples/prototypes/test_alloc.mlir
A mlir/examples/prototypes/test_gpu.mlir
A mlir/examples/prototypes/test_index.mlir
A mlir/examples/prototypes/test_rocdl.mlir
A mlir/examples/prototypes/test_vector.mlir
A mlir/examples/prototypes/toy_threadwise_copy_filter_tensor.mlir
A mlir/examples/prototypes/toy_threadwise_copy_input_tensor.mlir
Log Message:
-----------
XXX stash all experimental MLIR codes.
Commit: b54d51fdc76e60423c30e270a6d86dbd3144f6ba
https://github.com/llvm/llvm-project/commit/b54d51fdc76e60423c30e270a6d86dbd3144f6ba
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpenOps/Transform/ConvertMIOpenOpsToStd.cpp
Log Message:
-----------
XXX tame lowering down to LLVM.
Commit: 01e15b94d9b6d650bb06e5b5deda151eeace836a
https://github.com/llvm/llvm-project/commit/01e15b94d9b6d650bb06e5b5deda151eeace836a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
A mlir/include/mlir/Conversion/MIOpenToGPU/MIOpenToGPU.h
M mlir/include/mlir/Conversion/Passes.td
M mlir/include/mlir/Dialect/CMakeLists.txt
A mlir/include/mlir/Dialect/MIOpen/CMakeLists.txt
A mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
A mlir/include/mlir/Dialect/MIOpen/MIOpenOps.h
A mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
A mlir/include/mlir/Dialect/MIOpen/Passes.h
A mlir/include/mlir/Dialect/MIOpen/Passes.td
R mlir/include/mlir/Dialect/MIOpenOps/CMakeLists.txt
R mlir/include/mlir/Dialect/MIOpenOps/ConvertMIOpenOpsToStd.h
R mlir/include/mlir/Dialect/MIOpenOps/LowerMIOpenOps.h
R mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
R mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h
R mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
R mlir/include/mlir/Dialect/MIOpenOps/Passes.h
M mlir/include/mlir/InitAllDialects.h
M mlir/include/mlir/InitAllPasses.h
M mlir/include/mlir/InitAllTranslations.h
A mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Conversion/CMakeLists.txt
A mlir/lib/Conversion/MIOpenToGPU/CMakeLists.txt
A mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/CMakeLists.txt
A mlir/lib/Dialect/MIOpen/CMakeLists.txt
A mlir/lib/Dialect/MIOpen/Driver/CMakeLists.txt
A mlir/lib/Dialect/MIOpen/Driver/mlir-miopen-driver.cpp
A mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
A mlir/lib/Dialect/MIOpen/Transforms/AffineTransforms.cpp
A mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
A mlir/lib/Dialect/MIOpen/Transforms/CMakeLists.txt
A mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
A mlir/lib/Dialect/MIOpen/Transforms/PassDetail.h
R mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
R mlir/lib/Dialect/MIOpenOps/CppOutput/CMakeLists.txt
R mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
R mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
R mlir/lib/Dialect/MIOpenOps/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
R mlir/lib/Dialect/MIOpenOps/DialectRegistration.cpp
R mlir/lib/Dialect/MIOpenOps/Driver/CMakeLists.txt
R mlir/lib/Dialect/MIOpenOps/Driver/mlir-miopen-driver.cpp
R mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
R mlir/lib/Dialect/MIOpenOps/MIOpenOps.cpp
R mlir/lib/Dialect/MIOpenOps/Transform/AffineTransforms.cpp
R mlir/lib/Dialect/MIOpenOps/Transform/AffixTuningParameters.cpp
R mlir/lib/Dialect/MIOpenOps/Transform/CMakeLists.txt
R mlir/lib/Dialect/MIOpenOps/Transform/ConvertMIOpenOpsToStd.cpp
M mlir/lib/Target/CMakeLists.txt
A mlir/lib/Target/CppOutput/CMakeLists.txt
A mlir/lib/Target/CppOutput/ConvertToMIOpenCPP.cpp
A mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
A mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
M mlir/lib/Transforms/CMakeLists.txt
Log Message:
-----------
Rebase to LLVM/MLIR tip.
Commit: 80ec9f069010a09ee28b0900ce08b86daed9dba8
https://github.com/llvm/llvm-project/commit/80ec9f069010a09ee28b0900ce08b86daed9dba8
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/examples/prototypes/prototype_threadwise_copy.mlir
A mlir/examples/prototypes/test_vector_reshape.mlir
A mlir/examples/prototypes/test_vector_transfer_loop.mlir
M mlir/examples/prototypes/toy_threadwise_copy_filter_tensor.mlir
Log Message:
-----------
XXX stash updated MLIR prototypes after rebase.
Commit: bc9e6e5cab73ba474fe256cc7410a83db472b35f
https://github.com/llvm/llvm-project/commit/bc9e6e5cab73ba474fe256cc7410a83db472b35f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
Log Message:
-----------
XXX WIP lowering MIOpen threadwise ops to various MLIR dialects.
Commit: 0fc520849f41a331ab35eb8ce5f6f88bd012e4e9
https://github.com/llvm/llvm-project/commit/0fc520849f41a331ab35eb8ce5f6f88bd012e4e9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Migrate miopen.subview lowering logic from -convert-miopen-to-gpu to -miopen-lowering-step3.
Commit: 5daf41b6dd54db6127fd4c037c93923b1613d0f8
https://github.com/llvm/llvm-project/commit/5daf41b6dd54db6127fd4c037c93923b1613d0f8
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Lowering logic for miopen.transform.
Commit: 37f9fa1dacb2c9205409d11ec04e55d0adbb6fa8
https://github.com/llvm/llvm-project/commit/37f9fa1dacb2c9205409d11ec04e55d0adbb6fa8
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
Log Message:
-----------
Revise miopen.gpu_alloc lowering logic after rebase.
Commit: abe187eb7705647e91003b310a9d51610b1ae744
https://github.com/llvm/llvm-project/commit/abe187eb7705647e91003b310a9d51610b1ae744
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Revise miopen.gridwise_gemm lowering logic for miopen.subview to carry affine maps on the result memref types.
Commit: 061793a45cc1c0677997a015d07e8ba443aa7082
https://github.com/llvm/llvm-project/commit/061793a45cc1c0677997a015d07e8ba443aa7082
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Revise miopen.subview lowering logic.
Pass the affine map to its users.
Commit: fe4b31f878b1ddeae2e52d795bd8141c589d7b34
https://github.com/llvm/llvm-project/commit/fe4b31f878b1ddeae2e52d795bd8141c589d7b34
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
Log Message:
-----------
Revise syntax of miopen.gpu_alloc().
Now miopen.gpu_alloc() would return memrefs of arbitrary ranks.
Commit: f178145d8535189ba0e161856ff0219779f002c3
https://github.com/llvm/llvm-project/commit/f178145d8535189ba0e161856ff0219779f002c3
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
Log Message:
-----------
Revise miopen.fill and how we allocate VGPR (address space 5).
Commit: bbfbd341c5f61cf9672dddfe1ee04d2984adebe0
https://github.com/llvm/llvm-project/commit/bbfbd341c5f61cf9672dddfe1ee04d2984adebe0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Change some VGPR allocations become 2D. Revise miopen.fill lowering logic.
Commit: 73799fcc595a14e985fa6624b84fc94885308f4b
https://github.com/llvm/llvm-project/commit/73799fcc595a14e985fa6624b84fc94885308f4b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Change how VGPRs are allocated.
Commit: 0dcee284088a9719ea1da6aa4329d56fd886eeee
https://github.com/llvm/llvm-project/commit/0dcee284088a9719ea1da6aa4329d56fd886eeee
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
minor refactoring.
Commit: 562edbca17ca28ab29f4a43829e52976b5c3ac4c
https://github.com/llvm/llvm-project/commit/562edbca17ca28ab29f4a43829e52976b5c3ac4c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Revise -miopen-lowering-step2 wrt gridwise_gemm lowering.
Now LDS buffers are passed to blockwise_gemm and blockwise_copy in 2D forms.
Commit: aadb30598da9742822d48be5690559113ec9d0b2
https://github.com/llvm/llvm-project/commit/aadb30598da9742822d48be5690559113ec9d0b2
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
Log Message:
-----------
Revise miopen.fill so it supports integer types. Add 3 2D coordinates for BlockwiseCopyOp.
Commit: 555dad6726c366667647662670cee6e31511b176
https://github.com/llvm/llvm-project/commit/555dad6726c366667647662670cee6e31511b176
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
Log Message:
-----------
Make memref types accepted by blockwise_gemm and threadwise_gemm strict.
Commit: 68f7555b0b968dbff5e5ba90b57dae2ab73399e0
https://github.com/llvm/llvm-project/commit/68f7555b0b968dbff5e5ba90b57dae2ab73399e0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
Log Message:
-----------
WIP revise BlockwiseCopyOp so it takes source / dest coordinates.
TBD add MovePos op.
TBD revise how "move_source_slice_window" is implemented.
Commit: a1f2936d75bfb9b47e025930a2201788aff1b3d3
https://github.com/llvm/llvm-project/commit/a1f2936d75bfb9b47e025930a2201788aff1b3d3
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Fix unit tests.
Commit: 7e6fa448d328d925b49e022c7d6c3a23722afe4d
https://github.com/llvm/llvm-project/commit/7e6fa448d328d925b49e022c7d6c3a23722afe4d
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Introduce miopen.move_pos operator.
Commit: da892359462bcb65f98e9a7e4a0a9ecf439d5cf9
https://github.com/llvm/llvm-project/commit/da892359462bcb65f98e9a7e4a0a9ecf439d5cf9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Use miopen.move_pos op in miopen.gridwise_gemm lowering.
Commit: 2c3d73715fa8006b3ebd12f03c2204c76d81fd84
https://github.com/llvm/llvm-project/commit/2c3d73715fa8006b3ebd12f03c2204c76d81fd84
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Implement lowering logic for miopen.move_pos to std.
Commit: 5193b911149b19e13eeece195557726caf6d670c
https://github.com/llvm/llvm-project/commit/5193b911149b19e13eeece195557726caf6d670c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Rename variables.
Commit: 228a600d6cb643b00a8423a3f1a594e84f417daf
https://github.com/llvm/llvm-project/commit/228a600d6cb643b00a8423a3f1a594e84f417daf
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Remove unused codes for move_source_slice_window.
Commit: cea251332a2232d8504122b298c3664ffa082d2b
https://github.com/llvm/llvm-project/commit/cea251332a2232d8504122b298c3664ffa082d2b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Revise stack allocations for src / dest coordinates for blockwise_copy for Matrix A and B.
Commit: 1ccf623592b852b591c02d3be92f8d9b00b11d69
https://github.com/llvm/llvm-project/commit/1ccf623592b852b591c02d3be92f8d9b00b11d69
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Revise miopen.threadwise_copy interface and revise its call sites accordingly.
- source and destination memrefs can be on any address spaces
- source and destination memrefs can carry affine maps which matches its rank. for naive tensors on LDS that's the case.
- source and destination memrefs can be addressed with externally specified affine maps via attributes. for generic tensors on global VRAM that's the case.
for example:
- for matrix A (filter) / matrix B (input) we can use 2D -> 4D transformations
- for matrix C (output) with cluster (m0, m1, n0, n1) we can use 4D -> 2D -> 4D transformations
- coordinate transformation maps in each stage thru miopen.transform could be kept. so we can decide whether to store coordinates during all these transformations. or just keep the most upper index.
Commit: ba0a41935faeac83cb4789fd40ed302e7aab4057
https://github.com/llvm/llvm-project/commit/ba0a41935faeac83cb4789fd40ed302e7aab4057
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Revise miopen.subview lowering logic to adopt new interface of miopen.threadwise_copy.
Disable miopen.transform lowering logic for now. Will fix it in the next commit.
Commit: c3665edad0f6331cfca500804b3599142b798f84
https://github.com/llvm/llvm-project/commit/c3665edad0f6331cfca500804b3599142b798f84
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Revise miopen.transform lowering logic to adopt new interface of miopen.threadwise_copy.
Commit: d7054ed041aed5ff4bef2f71ab3f51c21d3fb272
https://github.com/llvm/llvm-project/commit/d7054ed041aed5ff4bef2f71ab3f51c21d3fb272
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
A mlir/examples/prototypes/threadwise_gemm.mlir
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Adding threadwisegemm lowering implementation
Commit: bec227db19532e6aa19ee94734443401b961f7bb
https://github.com/llvm/llvm-project/commit/bec227db19532e6aa19ee94734443401b961f7bb
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Removing funcop and callop associated with threadwisegemm
Commit: 5e336db89f2a204620847cebe1367546fd75eda0
https://github.com/llvm/llvm-project/commit/5e336db89f2a204620847cebe1367546fd75eda0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/Passes.h
M mlir/include/mlir/Dialect/MIOpen/Passes.td
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/Transforms/CMakeLists.txt
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Introduce -miopen-lowering-step4 pass.
Commit: 40dd6df424e93ab12bdbb2cc7d3eb9def71fcde7
https://github.com/llvm/llvm-project/commit/40dd6df424e93ab12bdbb2cc7d3eb9def71fcde7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Skeleton of miopen.threadwise_copy lowering logic.
Move from -convert-miopen-to-gpu to -miopen-lowering-step4.
Commit: 2d16a35f4f8c03fa0427815044785b198ebacc0c
https://github.com/llvm/llvm-project/commit/2d16a35f4f8c03fa0427815044785b198ebacc0c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
Log Message:
-----------
Minor refactoring.
Commit: 7dc132a1691a6b78ca7af6c174bfb51655c5d42f
https://github.com/llvm/llvm-project/commit/7dc132a1691a6b78ca7af6c174bfb51655c5d42f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/examples/prototypes/test_vector_reshape.mlir
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Skeleton miopen.threadwise_copy prototype. Just make sure emitting vector dialect works.
Commit: bb342844d860bbce7acac15409ce29720947d066
https://github.com/llvm/llvm-project/commit/bb342844d860bbce7acac15409ce29720947d066
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
miopen.lds_barrier -> miopen.workgroup_barrier.
Disassociate hardware-specific nomenclature.
Commit: a561d84cbd0e5db9a4a8766d4b2cba748e9ee714
https://github.com/llvm/llvm-project/commit/a561d84cbd0e5db9a4a8766d4b2cba748e9ee714
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Add miopen.workgroup_id and miopen.workitem_id ops.
These ops would be used in grdwise_gemm, blockwise_gemm lowering steps.
Commit: 38360c21fc5bb5fa9b0231a5f8e253f81f8d35e2
https://github.com/llvm/llvm-project/commit/38360c21fc5bb5fa9b0231a5f8e253f81f8d35e2
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
Log Message:
-----------
Add lowering logic from miopen to GPU for indexing.
Commit: 2aa52066dddac0db6142cd4b50cfc1a7a775ae0f
https://github.com/llvm/llvm-project/commit/2aa52066dddac0db6142cd4b50cfc1a7a775ae0f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Rename gpu_alloc -> alloc.
Commit: b2b1b4648749cdd0225282a8b32ed51e899154d0
https://github.com/llvm/llvm-project/commit/b2b1b4648749cdd0225282a8b32ed51e899154d0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Skeleton of populating attributes for blockwise_copy from information gathered in gridwise_gemm.
Commit: 244025a8199ef39aa4dfaf25f283b5300b3e5136
https://github.com/llvm/llvm-project/commit/244025a8199ef39aa4dfaf25f283b5300b3e5136
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
Add comments on the usage of all template arguments used by gridwise_gemm and how they can be computed from attributes.
Commit: 7884b36a9e335b3f22265d7e026705682ca575da
https://github.com/llvm/llvm-project/commit/7884b36a9e335b3f22265d7e026705682ca575da
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Add comments on the usage of all template arguments used by blockwise_copy and how they can be computed from attributes.
Commit: 0d18d9c98eb5f8a8fa3d231956a7f237583624dc
https://github.com/llvm/llvm-project/commit/0d18d9c98eb5f8a8fa3d231956a7f237583624dc
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Add comments on the usage of all template arguments used by blockwise_gemm and how they can be computed from attributes.
Commit: 57e03e762bb99ec085abf1c798b9983e772af02e
https://github.com/llvm/llvm-project/commit/57e03e762bb99ec085abf1c798b9983e772af02e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Compute m_block_data_on_global and n_block_data_on_global and pass to blockwise_copy statements.
Commit: 8043c84bfc424183f5180b1f0b57fa6b4da89f9a
https://github.com/llvm/llvm-project/commit/8043c84bfc424183f5180b1f0b57fa6b4da89f9a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Revise comments.
Commit: 40a80021af4158b1e7fe03b14c38b3afbdf178fe
https://github.com/llvm/llvm-project/commit/40a80021af4158b1e7fe03b14c38b3afbdf178fe
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
Log Message:
-----------
Set kernel attribute for each GPUFuncOp.
Notice we might want to revisit this later on.
Commit: a9a761cc429884ce2d99d054b7372cbefb62fa2e
https://github.com/llvm/llvm-project/commit/a9a761cc429884ce2d99d054b7372cbefb62fa2e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
Log Message:
-----------
Remove FuncOp in -convert-miopen-to-gpu conversion pass.
Commit: ec87ba61028e4a57d472be2caef984c2df5347e4
https://github.com/llvm/llvm-project/commit/ec87ba61028e4a57d472be2caef984c2df5347e4
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp
Log Message:
-----------
Improve -mlir-to-rocdlir so it could translate a GPU module when it sees one.
Commit: fb33bfa7c0b51345ba457cf3a080037a11d4ee37
https://github.com/llvm/llvm-project/commit/fb33bfa7c0b51345ba457cf3a080037a11d4ee37
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Skeleton logic assigning attributes to threadwise_copy.
Two improvements needed:
1) distinguish between copies between naive tensors and generic tensors.
2) translate matrix A/B/C to source / dest memrefs.
Commit: 659150139311d68fa1b911e7484e2d13d47c6d7c
https://github.com/llvm/llvm-project/commit/659150139311d68fa1b911e7484e2d13d47c6d7c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Improve blockwise_gemm lowering logic.
Commit: f1b53dfcdd28f49a90ec898bc7a77f5615ca5887
https://github.com/llvm/llvm-project/commit/f1b53dfcdd28f49a90ec898bc7a77f5615ca5887
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
XXX. WIP. Revise comments and logic on affixing attributes for threadwise_copy ops.
NOT completed yet.
Commit: fc4435a89fa91c623394d30a00eaa29e7961cdd5
https://github.com/llvm/llvm-project/commit/fc4435a89fa91c623394d30a00eaa29e7961cdd5
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Compute cluster lengths and thread slice lengths for blockwise_copy calls within gridwise_gemm lowering.
Commit: 6a10db39540819bde08cf97312ed7b7e35cae9b9
https://github.com/llvm/llvm-project/commit/6a10db39540819bde08cf97312ed7b7e35cae9b9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Add coordinate computation logic in BlockwiseCopy ctor.
Commit: 96ba7a7ad92d78bd9bc544c3b3eaa2614caa21dc
https://github.com/llvm/llvm-project/commit/96ba7a7ad92d78bd9bc544c3b3eaa2614caa21dc
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
Log Message:
-----------
Revise blockwise_copy syntax to allow staging buffer.
Use the staging buffer when copying from global to LDS.
Compute the coordinates in blockwise_copy lowering logic.
Commit: 18798958bd1de7ac34950ac010f4e8c7ab388f65
https://github.com/llvm/llvm-project/commit/18798958bd1de7ac34950ac010f4e8c7ab388f65
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Fix one missing attribute from gridwise_gemm -> blockwise_gemm.
Commit: a89becdb18d4c8c4deda44a9b18e711914821347
https://github.com/llvm/llvm-project/commit/a89becdb18d4c8c4deda44a9b18e711914821347
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Improving attribute affixing logic.
- gridwise_gemm -> blockwise_copy
- gridwise_gemm -> blockwise_gemm
- gridwise_gemm -> threadwise_copy
- blockwise_copy -> threadwise_copy
- blockwise_gemm -> threadwise_copy
Commit: 3fe533ba9762f7571d09ff2fcf858972c671ac64
https://github.com/llvm/llvm-project/commit/3fe533ba9762f7571d09ff2fcf858972c671ac64
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Revise syntax of blockwise_gemm so it can take threadwise offset for Matrix A and B.
Commit: 9683daf3ec8f047af333b9cb4197096dd5d1bc51
https://github.com/llvm/llvm-project/commit/9683daf3ec8f047af333b9cb4197096dd5d1bc51
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Compute source and dest coordinates in blockwise_gemm -> threadwise_copy.
Commit: 292108f5218fc3a625d606b3f380e735006b19c2
https://github.com/llvm/llvm-project/commit/292108f5218fc3a625d606b3f380e735006b19c2
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Reduce calls to getOperand(), use dialect-specific parameter names instead.
Commit: d8c09f87ce836e65ddc4c08b0c9189be07fac83d
https://github.com/llvm/llvm-project/commit/d8c09f87ce836e65ddc4c08b0c9189be07fac83d
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Compute threadwise_copy coordinates used in gridwise_gemm.
Commit: 00e0c2812e65c8d19c7f8ddd7ae0289ffb887be5
https://github.com/llvm/llvm-project/commit/00e0c2812e65c8d19c7f8ddd7ae0289ffb887be5
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Improve comments in blockwise_copy -> threadwise_copy attribute affixing logic.
Commit: 5d027d17fbbe0b4191144054ea3bf40e2f6c8ff3
https://github.com/llvm/llvm-project/commit/5d027d17fbbe0b4191144054ea3bf40e2f6c8ff3
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
Change tuning parameters and attributes to cope that we are only using 2D coordinates storing VGPR to global.
Commit: 0e410c6979d045c171af84f3805e9b275b9d3227
https://github.com/llvm/llvm-project/commit/0e410c6979d045c171af84f3805e9b275b9d3227
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Refactor uses of op.getLoc().
Commit: 9488261ab4438060a467f03147c1358b56aeb553
https://github.com/llvm/llvm-project/commit/9488261ab4438060a467f03147c1358b56aeb553
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
First implementation of threeadwise_copy when source and dest are both naive tensors.
This commit realizes ThreadwiseSliceCopy.
Commit: 45d62ccf38f20c0f108199c4333d2e3eadd45cd6
https://github.com/llvm/llvm-project/commit/45d62ccf38f20c0f108199c4333d2e3eadd45cd6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/Passes.h
M mlir/include/mlir/Dialect/MIOpen/Passes.td
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Introduce -miopen-lowering-step5 pass.
Commit: 014b84f4ecbf761b2d2266a85f986424d96d3a8c
https://github.com/llvm/llvm-project/commit/014b84f4ecbf761b2d2266a85f986424d96d3a8c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
First implementation of threeadwise_copy considering generic tensors.
Commit: cdea5aa31e589b18a32f60af73256900fdbd88f6
https://github.com/llvm/llvm-project/commit/cdea5aa31e589b18a32f60af73256900fdbd88f6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/CMakeLists.txt
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/InitAllDialects.h
M mlir/include/mlir/InitAllPasses.h
M mlir/lib/Conversion/CMakeLists.txt
M mlir/lib/Dialect/CMakeLists.txt
M mlir/lib/Dialect/MIOpen/Transforms/CMakeLists.txt
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
M mlir/lib/Target/CppOutput/CMakeLists.txt
M mlir/lib/Transforms/CMakeLists.txt
Log Message:
-----------
Rebase to tip of LLVM / MLIR on 5/18/2020.
- Loop dialect is now renamed as SCF.
Commit: 7e29f5e6adadba0151bb1534dcecfd38e6c74a35
https://github.com/llvm/llvm-project/commit/7e29f5e6adadba0151bb1534dcecfd38e6c74a35
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp
Log Message:
-----------
Improve logic in translateToROCDLIR so it works on either ModuleOp or GPUModuleOp.
Commit: 09e3965ae6dec0468b52b4a48af94fc7c0727bc9
https://github.com/llvm/llvm-project/commit/09e3965ae6dec0468b52b4a48af94fc7c0727bc9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Target/CppOutput/CMakeLists.txt
Log Message:
-----------
Rebase to tip of LLVM / MLIR on 5/28/2020.
Commit: 8cbd5aebcc63fcf5d852e99a7043cab72ad2e507
https://github.com/llvm/llvm-project/commit/8cbd5aebcc63fcf5d852e99a7043cab72ad2e507
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
M mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h
M mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp
M mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
M mlir/test/Target/rocdl.mlir
Log Message:
-----------
[mlir][llvm] allow mlir-translate carry custom triple and data layout.
Summary:
- Teach mlir-translate to use custom triple and data layout.
- Change convert-to-rocdlir pass to pass AMDGPU-specific triple and target
layout string.
- Amend test case to check alloca on non-zero addrspace.
Reviewers: ftynse, nicolasvasilache, mehdi_amini
Subscribers: tpr, mehdi_amini, rriddle, jpienaar, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, grosul1, frgossen, Kayjukh, llvm-commits
Tags: #llvm #mlir
Differential Revision: https://reviews.llvm.org/D79019
Commit: 1b15b2b0b4765c604053b3bb0819eecb9810237b
https://github.com/llvm/llvm-project/commit/1b15b2b0b4765c604053b3bb0819eecb9810237b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/GPU/GPUOps.td
M mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
M mlir/test/Dialect/GPU/invalid.mlir
M mlir/test/Dialect/GPU/ops.mlir
Log Message:
-----------
[mlir][gpu] Fix logic error in D79508 computing number of private attributions.
Fix logic error in D79508. The old logic would make the first check in
`GPUFuncOp::verifyBody` always pass.
Commit: 3580313d44833ac4f80e85b554ad829629e7d962
https://github.com/llvm/llvm-project/commit/3580313d44833ac4f80e85b554ad829629e7d962
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
A mlir/test/mlir-rocm-runner/vecadd_transfer.mlir
M mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp
Log Message:
-----------
Add an example with explicit data transfers between CPU and GPU.
Commit: 2fc0e8726f0d75df14eca1429436ff372960fa62
https://github.com/llvm/llvm-project/commit/2fc0e8726f0d75df14eca1429436ff372960fa62
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Driver/CMakeLists.txt
M mlir/lib/Dialect/MIOpen/Driver/mlir-miopen-driver.cpp
Log Message:
-----------
Modernize mlir-miopen-driver logic.
Commit: 8159cb5ded6ce8ab29c457b49264241533c3c703
https://github.com/llvm/llvm-project/commit/8159cb5ded6ce8ab29c457b49264241533c3c703
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/CMakeLists.txt
M mlir/test/CMakeLists.txt
A mlir/test/Dialect/MIOpen/Driver/lit.local.cfg
A mlir/test/Dialect/MIOpen/Driver/populate.mlir
M mlir/test/lit.site.cfg.py.in
Log Message:
-----------
Add config option and test for mlir-miopen-driver.
-DMLIR_MIOPEN_DRIVER=1
Commit: 5283338699e11fc7b10a1658de5c4719b671c91a
https://github.com/llvm/llvm-project/commit/5283338699e11fc7b10a1658de5c4719b671c91a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Driver/mlir-miopen-driver.cpp
Log Message:
-----------
Refactor module population logic.
Commit: ed1c0ffeb6fd084e08be4109522df4e59d345b1d
https://github.com/llvm/llvm-project/commit/ed1c0ffeb6fd084e08be4109522df4e59d345b1d
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/CMakeLists.txt
R mlir/lib/Dialect/MIOpen/Driver/CMakeLists.txt
R mlir/lib/Dialect/MIOpen/Driver/mlir-miopen-driver.cpp
M mlir/tools/CMakeLists.txt
A mlir/tools/mlir-miopen-driver/CMakeLists.txt
A mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Move mlir-miopen-driver to its own directory under mlir/tools.
Commit: ec9e3b711bfc4c9367009575d384c3ab27a3d1c6
https://github.com/llvm/llvm-project/commit/ec9e3b711bfc4c9367009575d384c3ab27a3d1c6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Let mlir-miopen-driver to recognize mlir-opt options.
Commit: 60b4cad11d3dbb9e83361dcef0add3de1d0f050a
https://github.com/llvm/llvm-project/commit/60b4cad11d3dbb9e83361dcef0add3de1d0f050a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Enable MIOpen dialect lowering pipeline in miopen-mlir-driver.
Commit: 394a48a097b3ab0cbac00698a292870a17472582
https://github.com/llvm/llvm-project/commit/394a48a097b3ab0cbac00698a292870a17472582
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Enable setting a host harness file.
Lowering with a host harness file specified in WIP.
Commit: c572f95c98109457bcf0a220e778865d751991c9
https://github.com/llvm/llvm-project/commit/c572f95c98109457bcf0a220e778865d751991c9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
R mlir/test/Dialect/MIOpen/Driver/lit.local.cfg
R mlir/test/Dialect/MIOpen/Driver/populate.mlir
A mlir/test/mlir-miopen-driver/lit.local.cfg
A mlir/test/mlir-miopen-driver/populate.mlir
Log Message:
-----------
Move mlir-miopen-driver test to its own directory.
Commit: 0a5f255ecd8562ca994f6d7e6968d9d40e176643
https://github.com/llvm/llvm-project/commit/0a5f255ecd8562ca994f6d7e6968d9d40e176643
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
A mlir/test/mlir-miopen-driver/conv2d_harness.mlir
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Introduce the concept of host harness to mlir-miopen-driver with a test case.
When `--host` is specified along with a filename, `mlir-miopen-driver` would
load it, parse it, and populate 2D convolution logic inside the entry point
function.
An entry point function could be specified with `--entry-point` flag, the
default value is `conv2d`. A `gpu.launch` is expected inside the entry point
function and 2D convolution logic would be populated inside the operation.
An example, `conv2d_harness.mlir` is included in this commit to depict the
usage.
The purpose of this commit is to allow quick prototyping and avoid
time-consuming logic to populate host-side logic.
Commit: f594a6f2bdb5512363e6a1d35455214813db268e
https://github.com/llvm/llvm-project/commit/f594a6f2bdb5512363e6a1d35455214813db268e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
A mlir/test/mlir-rocm-runner/vecadd_transfer_launch_func.mlir
Log Message:
-----------
Introduce one unit test which uses gpu.launch_func and workgroup/private attributions.
Commit: 6caf7288555ad6c35d7be8c51a303688dca18378
https://github.com/llvm/llvm-project/commit/6caf7288555ad6c35d7be8c51a303688dca18378
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/test/mlir-miopen-driver/conv2d_harness.mlir
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Revise how host harness is used in mlir-miopen-driver.
Instead using `gpu.launc`, use `gpu.launch_func` instead so workgroup and
private attributions of GPU dialect can be used properly.
Commit: ec53e851893d1af6466889fc0cdbb91b81f61a28
https://github.com/llvm/llvm-project/commit/ec53e851893d1af6466889fc0cdbb91b81f61a28
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Conversion/MIOpenToGPU/MIOpenToGPU.h
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/test/mlir-miopen-driver/conv2d_harness.mlir
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Improve test case for mlir-miopen-driver to cover lowering process.
Commit: 1ef6153e6ce10ba102163a01e1fd6533673784e1
https://github.com/llvm/llvm-project/commit/1ef6153e6ce10ba102163a01e1fd6533673784e1
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Let mlir-miopen-driver to populate default values always.
Commit: 93defa17845dd5edf7a7a5888f1326b90edbc1e1
https://github.com/llvm/llvm-project/commit/93defa17845dd5edf7a7a5888f1326b90edbc1e1
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp
Log Message:
-----------
HIP runtime wrappers for 4D f32 memrefs.
Commit: b2cb95ec526fe35299bb37ff8d6bdc932ec88171
https://github.com/llvm/llvm-project/commit/b2cb95ec526fe35299bb37ff8d6bdc932ec88171
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/test/mlir-miopen-driver/conv2d_harness.mlir
Log Message:
-----------
Add E2E check for mlir-miopen-driver.
TBD. Only check the content of filter tensor at this moment.
Commit: 00fd77663cf420d6a229ea21b3c16c257c7444c5
https://github.com/llvm/llvm-project/commit/00fd77663cf420d6a229ea21b3c16c257c7444c5
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops.mlir
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
add wrw parser
Commit: 8d73e836f2ccf8ec2d9a685660b20beeef31dc3e
https://github.com/llvm/llvm-project/commit/8d73e836f2ccf8ec2d9a685660b20beeef31dc3e
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
fixed
Commit: 55da2c47a1f8217cd739f18ab99ed5e0ef79878c
https://github.com/llvm/llvm-project/commit/55da2c47a1f8217cd739f18ab99ed5e0ef79878c
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
M mlir/test/Dialect/MIOpen/lowering_filter_tensor_ckyx_cnhw_knhw.mlir
M mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_position_cyxk_chwn_khwn.mlir
M mlir/test/Dialect/MIOpen/lowering_input_tensor_cyxk_cnhw_hnhw.mlir
M mlir/test/Dialect/MIOpen/lowering_memref_kcyx_nchw_nkhw.mlir
M mlir/test/Dialect/MIOpen/lowering_output_tensor_kyxc_nhwc_nhwk.mlir
M mlir/test/Dialect/MIOpen/lowering_top_level.mlir
M mlir/test/mlir-miopen-driver/conv2d_harness.mlir
M mlir/test/mlir-rocm-runner/gpu-to-hsaco.mlir
M mlir/test/mlir-rocm-runner/two-modules.mlir
M mlir/test/mlir-rocm-runner/vecadd.mlir
M mlir/test/mlir-rocm-runner/vecadd_transfer.mlir
M mlir/test/mlir-rocm-runner/vecadd_transfer_launch_func.mlir
Log Message:
-----------
add lowering
Commit: 95391d802bcfd1b237f7607f9ff9050e1ce84abd
https://github.com/llvm/llvm-project/commit/95391d802bcfd1b237f7607f9ff9050e1ce84abd
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/test/mlir-rocm-runner/gpu-to-hsaco.mlir
M mlir/test/mlir-rocm-runner/two-modules.mlir
M mlir/test/mlir-rocm-runner/vecadd.mlir
M mlir/test/mlir-rocm-runner/vecadd_transfer.mlir
M mlir/test/mlir-rocm-runner/vecadd_transfer_launch_func.mlir
Log Message:
-----------
remove patch from test
Commit: 84bf61833b5ae327418685a7d9c9b13571781dc8
https://github.com/llvm/llvm-project/commit/84bf61833b5ae327418685a7d9c9b13571781dc8
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/test/mlir-miopen-driver/conv2d_harness.mlir
Log Message:
-----------
remove patch from test
Commit: 496bd4e06d3141c5b98e31e22bc54f322e29d07b
https://github.com/llvm/llvm-project/commit/496bd4e06d3141c5b98e31e22bc54f322e29d07b
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-05 (Fri, 05 Jun 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
fixed func.getArgument
Commit: 4b28ac9fa64c8dc5da8edc401b27b0840641c374
https://github.com/llvm/llvm-project/commit/4b28ac9fa64c8dc5da8edc401b27b0840641c374
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-06 (Sat, 06 Jun 2020)
Changed paths:
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
add cpp translation
Commit: 6d2d553237366f8da93ba948ea8760646eef0f3c
https://github.com/llvm/llvm-project/commit/6d2d553237366f8da93ba948ea8760646eef0f3c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-08 (Mon, 08 Jun 2020)
Changed paths:
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Merge pull request #17 from zjing14/add_wrw_cpp
Add wrw cpp translation
Commit: e71660642427f718df48c69a789573a06798d95c
https://github.com/llvm/llvm-project/commit/e71660642427f718df48c69a789573a06798d95c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-08 (Mon, 08 Jun 2020)
Changed paths:
M mlir/test/mlir-miopen-driver/conv2d_harness.mlir
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Introduce blockSize and gridSize parameter to mlir-miopen-driver.
This allows more detailed debugging.
Also change the e2e test to check output tensor.
Commit: d5005f6f6cfc3dc11e3f0de1b684c55b94a33688
https://github.com/llvm/llvm-project/commit/d5005f6f6cfc3dc11e3f0de1b684c55b94a33688
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-08 (Mon, 08 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Use constants defined in GPU dialect.
Commit: b752611e6c6b58cbaa1e94dfef688d2f50a7fb60
https://github.com/llvm/llvm-project/commit/b752611e6c6b58cbaa1e94dfef688d2f50a7fb60
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-08 (Mon, 08 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Fix m_block_data_on_global and n_block_data_on_global logic.
Commit: 9cb586fa5ee3d63ac9ad4fb80afe664e685084bc
https://github.com/llvm/llvm-project/commit/9cb586fa5ee3d63ac9ad4fb80afe664e685084bc
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-08 (Mon, 08 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Fix logic to compute level0_m_id and level0_n_id.
Commit: c4d6b0b786c0b0ffa5515cce416c1b74a055d746
https://github.com/llvm/llvm-project/commit/c4d6b0b786c0b0ffa5515cce416c1b74a055d746
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-08 (Mon, 08 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Fix incorrect comment in miopen.blockwise_gemm lowering procedure.
Commit: 0d26b6ff9ba9b97b99b021920518c41673112124
https://github.com/llvm/llvm-project/commit/0d26b6ff9ba9b97b99b021920518c41673112124
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-08 (Mon, 08 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Re-indent comments.
Commit: 2ab02208fed80dadfc89e8bee3920198e1a1a569
https://github.com/llvm/llvm-project/commit/2ab02208fed80dadfc89e8bee3920198e1a1a569
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-09 (Tue, 09 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Re-indent comments. Identified one place within blockwise_gemm which might be faulty.
Commit: fcad5f40799ee5a3a6e7a9048fea6291ed202320
https://github.com/llvm/llvm-project/commit/fcad5f40799ee5a3a6e7a9048fea6291ed202320
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-09 (Tue, 09 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
Fix MPerLevel1Cluster, NPerLevel1Cluster, KPerThread.
After discussing with Chao, KPerThread should actually be 1, not 4.
Enable the fix to MPerLevel1Cluster and NPerLevel1Cluster identified in the
previous commit.
Now in 1 thread, 64 elements on VGPR can be computed instead of 16.
Commit: 87ec5d86a7dda6e47517b908236d0c123b1a4cc7
https://github.com/llvm/llvm-project/commit/87ec5d86a7dda6e47517b908236d0c123b1a4cc7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-09 (Tue, 09 Jun 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Revise the comment on the logic to compute gridSize in kernel launch logic.
Commit: 1210363510fbb22857ffd90a99fe8ff5d9ae9ae7
https://github.com/llvm/llvm-project/commit/1210363510fbb22857ffd90a99fe8ff5d9ae9ae7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-09 (Tue, 09 Jun 2020)
Changed paths:
M mlir/include/mlir/Conversion/MIOpenToGPU/MIOpenToGPU.h
M mlir/include/mlir/Conversion/Passes.td
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
A mlir/test/Conversion/MIOpenToGPU/alloc.mlir
A mlir/test/Conversion/MIOpenToGPU/emptykernel.mlir
A mlir/test/Conversion/MIOpenToGPU/misc_ops.mlir
Log Message:
-----------
Teach -convert-miopen-to-gpu to recognize custom kernel name. Add tests.
Commit: bd2a3f47a63f698829a6fc180272c8a0ad259e99
https://github.com/llvm/llvm-project/commit/bd2a3f47a63f698829a6fc180272c8a0ad259e99
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-09 (Tue, 09 Jun 2020)
Changed paths:
M mlir/include/mlir/Conversion/MIOpenToGPU/MIOpenToGPU.h
M mlir/include/mlir/Conversion/Passes.td
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/test/Conversion/MIOpenToGPU/emptykernel.mlir
Log Message:
-----------
Teach -convert-miopen-to-gpu to recognize custom GPU module name. Revise test.
Commit: 7bec2760d09f6ea4b5c906c1c88dae13859f1add
https://github.com/llvm/llvm-project/commit/7bec2760d09f6ea4b5c906c1c88dae13859f1add
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-09 (Tue, 09 Jun 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
A mlir/test/Conversion/MIOpenToGPU/existing_gpu_module.mlir
Log Message:
-----------
Teach -convert-miopen-to-gpu to use existing GPUModule if found.
Commit: 295b7380ed1ba9e139941b87f38d2e941de27189
https://github.com/llvm/llvm-project/commit/295b7380ed1ba9e139941b87f38d2e941de27189
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-09 (Tue, 09 Jun 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
A mlir/test/Conversion/MIOpenToGPU/existing_gpu_func.mlir
Log Message:
-----------
Teach -convert-miopen-to-gpu to use existing GPUFuncOp if found.
Commit: b1b8251bf4f9fc09c98053487092dc9e537006cf
https://github.com/llvm/llvm-project/commit/b1b8251bf4f9fc09c98053487092dc9e537006cf
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-10 (Wed, 10 Jun 2020)
Changed paths:
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
M mlir/test/Dialect/MIOpen/translate_bwd_data_kcyx_nchw_nkhw.mlir
M mlir/test/Dialect/MIOpen/translate_forward_kcyx_nchw_nkhw.mlir
Log Message:
-----------
Rewor tuning parameter generation
Commit: 3be64fc6f7f0ee2726914b47773d41ea79ed782b
https://github.com/llvm/llvm-project/commit/3be64fc6f7f0ee2726914b47773d41ea79ed782b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-10 (Wed, 10 Jun 2020)
Changed paths:
M mlir/include/mlir/Conversion/MIOpenToGPU/MIOpenToGPU.h
M mlir/include/mlir/Conversion/Passes.td
M mlir/include/mlir/Dialect/MIOpen/Passes.td
M mlir/lib/Conversion/MIOpenToGPU/CMakeLists.txt
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
A mlir/test/Conversion/MIOpenToGPU/e2e/lds_fillzero.mlir
Log Message:
-----------
Initial commit to allow fine-grained e2e tests for MIOpen dialect.
- Added -test-miopen-lowering-gpu-module conversion pass which operate on
gpu::GPUModuleOp instead of ModuleOp.
- Populated patterns in MIOpen dialect into -test-miopen-lowering-gpu-module.
- Create an e2e test, it uses:
- miopen.alloc
- miopen.fill
- miopen.workgroup_barrier
- miopen.workitem_id
Commit: fdea153874b9a4038fa943f67961bb43dd81412c
https://github.com/llvm/llvm-project/commit/fdea153874b9a4038fa943f67961bb43dd81412c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-10 (Wed, 10 Jun 2020)
Changed paths:
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Add a unit test for parsing of miopen.threadwise_gemm.
Commit: 7617e06c86e548b0467b41f2854345c65585c205
https://github.com/llvm/llvm-project/commit/7617e06c86e548b0467b41f2854345c65585c205
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-10 (Wed, 10 Jun 2020)
Changed paths:
A mlir/include/mlir/Dialect/MIOpen/gridwise_convolution_implicit_gemm_util.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
refactored shared functions into a header file
Commit: 797fec9527c883e1ae92c8c4c0247931f7fc50ef
https://github.com/llvm/llvm-project/commit/797fec9527c883e1ae92c8c4c0247931f7fc50ef
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-10 (Wed, 10 Jun 2020)
Changed paths:
A mlir/include/mlir/Dialect/MIOpen/gridwise_convolution_implicit_gemm_util.h
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
M mlir/test/Dialect/MIOpen/translate_bwd_data_kcyx_nchw_nkhw.mlir
M mlir/test/Dialect/MIOpen/translate_forward_kcyx_nchw_nkhw.mlir
Log Message:
-----------
Merge pull request #18 from jerryyin/miopen-dialect
Rework tuning parameter generation
Commit: 6a5525ee238826352d68ee4b420df3a8167f134e
https://github.com/llvm/llvm-project/commit/6a5525ee238826352d68ee4b420df3a8167f134e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-10 (Wed, 10 Jun 2020)
Changed paths:
A mlir/test/Conversion/MIOpenToGPU/e2e/threadwise_gemm.mlir
M mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp
Log Message:
-----------
Add an e2e test for miopen.threadwise_gemm.
Commit: 14c9daac94cd59a46c0c6278516240f9aa44288c
https://github.com/llvm/llvm-project/commit/14c9daac94cd59a46c0c6278516240f9aa44288c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-10 (Wed, 10 Jun 2020)
Changed paths:
A mlir/include/mlir/Dialect/MIOpen/gridwise_convolution_implicit_gemm_util.h
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
M mlir/test/Dialect/MIOpen/translate_bwd_data_kcyx_nchw_nkhw.mlir
M mlir/test/Dialect/MIOpen/translate_forward_kcyx_nchw_nkhw.mlir
Log Message:
-----------
Merge branch 'miopen-dialect' of github.com:whchung/llvm-project into miopen-dialect
Commit: c4c1fc80cdc1a6d352463a6ef7355433bb3ef010
https://github.com/llvm/llvm-project/commit/c4c1fc80cdc1a6d352463a6ef7355433bb3ef010
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-10 (Wed, 10 Jun 2020)
Changed paths:
A mlir/test/Conversion/MIOpenToGPU/e2e/lit.local.cfg
Log Message:
-----------
Add lit.local.cfg so e2e tests are not carried out if mlir-rocm-runner is not configured.
Commit: b316228cffbb846402de3b0fe0d920548f9e95a6
https://github.com/llvm/llvm-project/commit/b316228cffbb846402de3b0fe0d920548f9e95a6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-10 (Wed, 10 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Use scalar load / store transferring data from/to generic tensor for now.
Commit: b93f6ed4aba578d26464b51d6c4b63225b6adf8b
https://github.com/llvm/llvm-project/commit/b93f6ed4aba578d26464b51d6c4b63225b6adf8b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-10 (Wed, 10 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Use scalar load / store transferring data from/to naive tensor for now.
Commit: 6da52643c0cbe8000c18158aed1030ecf370d0e6
https://github.com/llvm/llvm-project/commit/6da52643c0cbe8000c18158aed1030ecf370d0e6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-15 (Mon, 15 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Fix Matrix A cluter / thread coordinate computation logic.
Commit: 2bd54749fda902ca103e059db2ac00b42c5762f2
https://github.com/llvm/llvm-project/commit/2bd54749fda902ca103e059db2ac00b42c5762f2
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-15 (Mon, 15 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Let miopen.threadwise_copy be aware of dim_access_order attribute.
Commit: 1fd9e9bae17bef8e0a5f03f5ad5dd823b9ad88e6
https://github.com/llvm/llvm-project/commit/1fd9e9bae17bef8e0a5f03f5ad5dd823b9ad88e6
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-16 (Tue, 16 Jun 2020)
Changed paths:
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
A mlir/test/Dialect/MIOpen/translate_cflags_bwd.mlir
A mlir/test/Dialect/MIOpen/translate_cflags_fw.mlir
Log Message:
-----------
Adding tuning parameters unit tests
* Populated function name for tuning parameter generation
* Added bwd and fw tests
Commit: e42ae14d0e00328cc6c46cd41074cf55df337f85
https://github.com/llvm/llvm-project/commit/e42ae14d0e00328cc6c46cd41074cf55df337f85
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-16 (Tue, 16 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Assorted changes to make matrix C VPGR -> global writeout to work.
- miopen.gridwise_gemm : add 4D -> 2D -> 4D writeout for matrix C.
- miopen.transform : fix a bug where attributes be dropped.
- miopen.transform : track the domain of upper-level coorindate.
- miopen.threadwise_copy : use domain attribute when both memrefs are generic.
Commit: 8c1697b810d5ab04708b55ac96d173c4860de442
https://github.com/llvm/llvm-project/commit/8c1697b810d5ab04708b55ac96d173c4860de442
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-16 (Tue, 16 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Fix coefficient for Embed.
Commit: 72a96a46f82222cbfebc0d93bcad82adba5e08f7
https://github.com/llvm/llvm-project/commit/72a96a46f82222cbfebc0d93bcad82adba5e08f7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-17 (Wed, 17 Jun 2020)
Changed paths:
A mlir/examples/prototypes/stage2.mlir
A mlir/examples/prototypes/stage3-hacked-a.mlir
A mlir/examples/prototypes/stage3-hacked-b.mlir
A mlir/examples/prototypes/stage3-hacked-check-lds-a.mlir
A mlir/examples/prototypes/stage3-hacked-check-lds.mlir
A mlir/examples/prototypes/stage3-hacked-check-vgpr-a.mlir
A mlir/examples/prototypes/stage3-hacked-check-vgpr.mlir
A mlir/examples/prototypes/stage3-hacked-data-is-there.mlir
A mlir/examples/prototypes/stage3-hacked-to-work.mlir
A mlir/examples/prototypes/stage3-hacked.mlir
A mlir/examples/prototypes/stage3.mlir
Log Message:
-----------
Add various intermediate IR during the debugging process.
Some of them could be reduced to unit tests.
Commit: dbe2a97754f6303cad8a276b722804d74ad433de
https://github.com/llvm/llvm-project/commit/dbe2a97754f6303cad8a276b722804d74ad433de
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-17 (Wed, 17 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
Remove obsolete comment.
Commit: 45e51460d3eb2161f01ee39f89294aab7b2c6573
https://github.com/llvm/llvm-project/commit/45e51460d3eb2161f01ee39f89294aab7b2c6573
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-17 (Wed, 17 Jun 2020)
Changed paths:
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
A mlir/test/Dialect/MIOpen/translate_cflags_bwd.mlir
A mlir/test/Dialect/MIOpen/translate_cflags_fw.mlir
Log Message:
-----------
Merge pull request #20 from jerryyin/miopen-dialect
Adding tuning parameters unit tests
Commit: d579a0cc1e9fd6950a2ee37c3c6351d256e950d6
https://github.com/llvm/llvm-project/commit/d579a0cc1e9fd6950a2ee37c3c6351d256e950d6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-17 (Wed, 17 Jun 2020)
Changed paths:
A mlir/utils/jenkins/Jenkinsfile.downstream
A mlir/utils/jenkins/Jenkinsfile.upstream
Log Message:
-----------
Add Jenkinsfile for upstream and downstream checks.
Commit: c8503c27efe89d25f5563bacaf0ef5f93c4a8221
https://github.com/llvm/llvm-project/commit/c8503c27efe89d25f5563bacaf0ef5f93c4a8221
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-18 (Thu, 18 Jun 2020)
Changed paths:
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Rework cpp trans
Commit: 5553233325a1c6c6e82df8da29185b3468ef5f6c
https://github.com/llvm/llvm-project/commit/5553233325a1c6c6e82df8da29185b3468ef5f6c
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-18 (Thu, 18 Jun 2020)
Changed paths:
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
add tuning
Commit: 0099bfb2e7180b77cde40dfa41672f8e294f1408
https://github.com/llvm/llvm-project/commit/0099bfb2e7180b77cde40dfa41672f8e294f1408
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-19 (Fri, 19 Jun 2020)
Changed paths:
A mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
Log Message:
-----------
add tuing test
Commit: 570f011f846b38bb4849ff3c6cb1081f547f5cf2
https://github.com/llvm/llvm-project/commit/570f011f846b38bb4849ff3c6cb1081f547f5cf2
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-18 (Thu, 18 Jun 2020)
Changed paths:
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
A mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
Log Message:
-----------
Merge pull request #22 from zjing14/wrw_tuning
Add Wrw tuning
Commit: be138a39e19cda89a2dd64ee15cedde097cae179
https://github.com/llvm/llvm-project/commit/be138a39e19cda89a2dd64ee15cedde097cae179
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-19 (Fri, 19 Jun 2020)
Changed paths:
M mlir/include/mlir/Target/MIOpenCPP.h
Log Message:
-----------
improve gemmB vector
Commit: 937dabf988a501d1a26535cd8a1929027ee4d7e1
https://github.com/llvm/llvm-project/commit/937dabf988a501d1a26535cd8a1929027ee4d7e1
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-19 (Fri, 19 Jun 2020)
Changed paths:
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/test/Dialect/MIOpen/translate_cflags_fw.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
Log Message:
-----------
add hw vector support
Commit: da12f3a825654b7fbe812fee384fe47ac4b914c7
https://github.com/llvm/llvm-project/commit/da12f3a825654b7fbe812fee384fe47ac4b914c7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-18 (Thu, 18 Jun 2020)
Changed paths:
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/test/Dialect/MIOpen/translate_cflags_fw.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
Log Message:
-----------
Merge pull request #23 from zjing14/wrw_gemmb_vector_improve
Improve vector load
Commit: b1a9a067229d54108c25e8c43643093e8d92423f
https://github.com/llvm/llvm-project/commit/b1a9a067229d54108c25e8c43643093e8d92423f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-18 (Thu, 18 Jun 2020)
Changed paths:
A mlir/test/mlir-miopen-driver/conv2d_harness_kyxc_nhwc_nhwk.mlir
Log Message:
-----------
Add a test based on NHWC layout.
Commit: 2ce6f466e9154b3411089d19070d6d74676bf3b5
https://github.com/llvm/llvm-project/commit/2ce6f466e9154b3411089d19070d6d74676bf3b5
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-19 (Fri, 19 Jun 2020)
Changed paths:
M mlir/include/mlir/Target/MIOpenCPP.h
Log Message:
-----------
fixed vecLen
Commit: a9117fd79ca73f1bba28863bf7d208b522fd1e82
https://github.com/llvm/llvm-project/commit/a9117fd79ca73f1bba28863bf7d208b522fd1e82
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-19 (Fri, 19 Jun 2020)
Changed paths:
M mlir/include/mlir/Target/MIOpenCPP.h
Log Message:
-----------
Merge pull request #24 from zjing14/tuning_fix
fixed vecLen in tuning
Commit: ab438a191d27a792ad430d2acc4b404c1427bb23
https://github.com/llvm/llvm-project/commit/ab438a191d27a792ad430d2acc4b404c1427bb23
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-06-19 (Fri, 19 Jun 2020)
Changed paths:
M mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
Log Message:
-----------
changed test
Commit: 7206595fb3287260fee046ee802318c30fec3d14
https://github.com/llvm/llvm-project/commit/7206595fb3287260fee046ee802318c30fec3d14
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-19 (Fri, 19 Jun 2020)
Changed paths:
M mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
Log Message:
-----------
Merge pull request #25 from zjing14/tuning_fix
Correct wrw tuning test
Commit: 9deb3f4f088761c6ed485e80e7301dc882405343
https://github.com/llvm/llvm-project/commit/9deb3f4f088761c6ed485e80e7301dc882405343
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-21 (Sun, 21 Jun 2020)
Changed paths:
A mlir/examples/prototypes/test_vector_matmul.mlir
Log Message:
-----------
Add one example of how threadwise_gemm may be implemented with vector dialect.
Commit: 9ebd09667a8db2eb07fa2b5b40dc6002887a65b6
https://github.com/llvm/llvm-project/commit/9ebd09667a8db2eb07fa2b5b40dc6002887a65b6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-21 (Sun, 21 Jun 2020)
Changed paths:
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
Log Message:
-----------
Merge branch 'miopen-dialect' of github.com:whchung/llvm-project into miopen-dialect
Commit: 2941938c7babb6fc6e59a7883591ad6176a00c3e
https://github.com/llvm/llvm-project/commit/2941938c7babb6fc6e59a7883591ad6176a00c3e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-21 (Sun, 21 Jun 2020)
Changed paths:
M mlir/examples/prototypes/test_vector_matmul.mlir
Log Message:
-----------
Use vector.flat_transpose instead of vector.transpose.
Commit: 145960c28b05649243192be8b4ff9c053fc55a11
https://github.com/llvm/llvm-project/commit/145960c28b05649243192be8b4ff9c053fc55a11
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-22 (Mon, 22 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/Passes.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Compute block size and grid size in -miopen-affix-params pass.
Adopt the sizes in mlir-miopen-driver.
Commit: 3c8c5bea61fd9e10910e73cfe9fdbe73ffe6a946
https://github.com/llvm/llvm-project/commit/3c8c5bea61fd9e10910e73cfe9fdbe73ffe6a946
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-23 (Tue, 23 Jun 2020)
Changed paths:
A mlir/test/mlir-miopen-driver/conv2d_harness_ckyx_cnhw_knhw.mlir
Log Message:
-----------
Add a test based on CNHW layout.
Commit: 1615a312a7043c85191b8f40b53875e3dd5a8ecd
https://github.com/llvm/llvm-project/commit/1615a312a7043c85191b8f40b53875e3dd5a8ecd
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-23 (Tue, 23 Jun 2020)
Changed paths:
A mlir/test/mlir-miopen-driver/conv2d_harness_cyxk_chwn_khwn.mlir
Log Message:
-----------
Add a test based on CHWN layout.
Commit: ad91a50780f79358a18720d0e52c8880a11c92be
https://github.com/llvm/llvm-project/commit/ad91a50780f79358a18720d0e52c8880a11c92be
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-23 (Tue, 23 Jun 2020)
Changed paths:
A mlir/test/mlir-miopen-driver/conv2d_harness_yxkc_hwnc_hwnk.mlir
Log Message:
-----------
Add a test based on HWNC layout.
Commit: ed357a495cb19e933e41ebcffb78f8c70b4a1492
https://github.com/llvm/llvm-project/commit/ed357a495cb19e933e41ebcffb78f8c70b4a1492
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-23 (Tue, 23 Jun 2020)
Changed paths:
A mlir/test/mlir-miopen-driver/conv2d_harness_yxck_hwcn_hwkn.mlir
Log Message:
-----------
Add a test based on HWCN layout.
Commit: 2797e5c028cdafd902cf6d3fab33b51bc1db8991
https://github.com/llvm/llvm-project/commit/2797e5c028cdafd902cf6d3fab33b51bc1db8991
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-23 (Tue, 23 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/gridwise_convolution_implicit_gemm_util.h
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Abstract tuning parameter logic to the implicit_gemm_utils
Commit: dcf3833093d07255aafb1a33c74169fda52114e4
https://github.com/llvm/llvm-project/commit/dcf3833093d07255aafb1a33c74169fda52114e4
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-23 (Tue, 23 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/gridwise_convolution_implicit_gemm_util.h
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/test/Dialect/MIOpen/translate_cflags_bwd.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_fw.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
Log Message:
-----------
Addressing review feedbacks
* Refactored parameter generation
* Updated unit tests
* Improved debug information and failure handling in affix pass
Commit: b4caed40b10bff3e08685e927b40aac63095c1be
https://github.com/llvm/llvm-project/commit/b4caed40b10bff3e08685e927b40aac63095c1be
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-23 (Tue, 23 Jun 2020)
Changed paths:
R mlir/include/mlir/Dialect/MIOpen/gridwise_convolution_implicit_gemm_util.h
A mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Renaming implicit gemm header
Commit: 1cfdc9ddc9d318dfc0ba5a2f8d8e5562edb9ddc5
https://github.com/llvm/llvm-project/commit/1cfdc9ddc9d318dfc0ba5a2f8d8e5562edb9ddc5
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-23 (Tue, 23 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/include/mlir/Target/MIOpenCPP.h
Log Message:
-----------
Decouple MIOpenCPP.h from tuning parameters
Commit: 2ca0d4eeda2b5e5edf1bd572f5f1ed6968809415
https://github.com/llvm/llvm-project/commit/2ca0d4eeda2b5e5edf1bd572f5f1ed6968809415
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-24 (Wed, 24 Jun 2020)
Changed paths:
M mlir/utils/jenkins/Jenkinsfile.downstream
Log Message:
-----------
Fix Jenkins to build mlir-miopen-driver.
Commit: fe73575cf85300e4f27b3843369c1a948196da5e
https://github.com/llvm/llvm-project/commit/fe73575cf85300e4f27b3843369c1a948196da5e
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-24 (Wed, 24 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
Fixing matrix C dest per write to 1 in affix params
Commit: bb08e033a30c049cd4aec83d94f6cf8f2f44f29c
https://github.com/llvm/llvm-project/commit/bb08e033a30c049cd4aec83d94f6cf8f2f44f29c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-24 (Wed, 24 Jun 2020)
Changed paths:
R mlir/include/mlir/Dialect/MIOpen/gridwise_convolution_implicit_gemm_util.h
A mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
M mlir/test/Dialect/MIOpen/translate_cflags_bwd.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_fw.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
Log Message:
-----------
Merge pull request #27 from jerryyin/miopen-dialect
Abstract tuning parameter logic to gridwise_gemm_params.h
Commit: 4a8d52aef1b56c1d8092802ac003264bf711a50b
https://github.com/llvm/llvm-project/commit/4a8d52aef1b56c1d8092802ac003264bf711a50b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-24 (Wed, 24 Jun 2020)
Changed paths:
A mlir/test/mlir-miopen-driver/populate_host.mlir
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Initial commit to populate host logic.
Commit: 8448cecb2f525a8ac4433d5040720b2df2267101
https://github.com/llvm/llvm-project/commit/8448cecb2f525a8ac4433d5040720b2df2267101
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-24 (Wed, 24 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/Passes.h
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Enable mlir-miopen-driver blockSize override
Commit: ffc7951f8c3bff58471bdc506ecf1bd32f199da0
https://github.com/llvm/llvm-project/commit/ffc7951f8c3bff58471bdc506ecf1bd32f199da0
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-06-24 (Wed, 24 Jun 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
Adding comments to explain blockSizeOverride
Commit: 800ddfab5e6df832af93df406973e60a74a56aa4
https://github.com/llvm/llvm-project/commit/800ddfab5e6df832af93df406973e60a74a56aa4
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-24 (Wed, 24 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/Passes.h
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Merge pull request #28 from jerryyin/miopen-dialect
Enable mlir-miopen-driver blockSize override
Commit: 43017b27f3b9d42a4ccc56466de9f5ba053313a7
https://github.com/llvm/llvm-project/commit/43017b27f3b9d42a4ccc56466de9f5ba053313a7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-24 (Wed, 24 Jun 2020)
Changed paths:
M mlir/test/mlir-miopen-driver/populate_host.mlir
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Improve mlir-miopen-driver to introduce host code generation logic.
-ph generate minimal host logic for testing.
Commit: bc75db9b90ab1c136f06ec178cfb244356f85d56
https://github.com/llvm/llvm-project/commit/bc75db9b90ab1c136f06ec178cfb244356f85d56
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-24 (Wed, 24 Jun 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/Passes.h
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Merge branch 'miopen-dialect' of github.com:whchung/llvm-project into miopen-dialect
Commit: 5c0cd3323304f995672a9fa8cb6917321421249b
https://github.com/llvm/llvm-project/commit/5c0cd3323304f995672a9fa8cb6917321421249b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-06-25 (Thu, 25 Jun 2020)
Changed paths:
A mlir/examples/prototypes/conv2d.mlir
Log Message:
-----------
Add one example lowering linald.conv on CPU.
Commit: bcbac645b7a199b9dd9af48ee8dfba6fa5c61b41
https://github.com/llvm/llvm-project/commit/bcbac645b7a199b9dd9af48ee8dfba6fa5c61b41
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-01 (Wed, 01 Jul 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Add passes to lower linalg dialect to CFG blocks.
Prepare for adding CPU verifier.
Commit: c8e43eaea43bbb14d6d430e21805d2efacfffbaf
https://github.com/llvm/llvm-project/commit/c8e43eaea43bbb14d6d430e21805d2efacfffbaf
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-01 (Wed, 01 Jul 2020)
Changed paths:
M mlir/examples/prototypes/conv2d.mlir
Log Message:
-----------
Improve lingalg.conv2d example with attributes.
Commit: f36fb543259e138eb28b9d1ba885a1941d4b54af
https://github.com/llvm/llvm-project/commit/f36fb543259e138eb28b9d1ba885a1941d4b54af
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-01 (Wed, 01 Jul 2020)
Changed paths:
M mlir/examples/prototypes/conv2d.mlir
Log Message:
-----------
Log the default layout used by linalg.conv .
Commit: 37597fdb16615c48ad674534a14005ae65433651
https://github.com/llvm/llvm-project/commit/37597fdb16615c48ad674534a14005ae65433651
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-01 (Wed, 01 Jul 2020)
Changed paths:
A mlir/test/mlir-miopen-driver/conv2d_harness_cpu_verifier.mlir
Log Message:
-----------
Add one end-to-end test with CPU verifier.
Commit: c445f8d9e1d26e2057131c5dec096864cf4b138c
https://github.com/llvm/llvm-project/commit/c445f8d9e1d26e2057131c5dec096864cf4b138c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-06 (Mon, 06 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Introduce miopen.mfma operator.
This operator would be lowered to gpu.mfma operator.
Commit: 15eaba31a1ec59e89309853549e751c328646dd6
https://github.com/llvm/llvm-project/commit/15eaba31a1ec59e89309853549e751c328646dd6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-06 (Mon, 06 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/GPU/GPUOps.td
M mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
M mlir/test/Dialect/GPU/ops.mlir
Log Message:
-----------
Introduce gpu.mfma operatior.
It would be lowered to one or more rocdl.mfma operators.
This operator is considered temporary and will be removed in the future when
miopen.mfma can be lowered to rocdl.mfma directly.
Commit: dee669b6a8bad8ccc7037fc2abaef1bed599c116
https://github.com/llvm/llvm-project/commit/dee669b6a8bad8ccc7037fc2abaef1bed599c116
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-06 (Mon, 06 Jul 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
A mlir/test/Conversion/MIOpenToGPU/mfma.mlir
Log Message:
-----------
Add conversion logic from miopen.mfma to gpu.mfma.
Commit: b44b95c89fcd3bbe33555d5794f325a31b1463fb
https://github.com/llvm/llvm-project/commit/b44b95c89fcd3bbe33555d5794f325a31b1463fb
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-06 (Mon, 06 Jul 2020)
Changed paths:
M mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
M mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
Log Message:
-----------
XXX. Experimental commit to lower gpu.mfma to rocdl.mfma.
Only consider very minimal case. The logic needs to be augmented to support
more general cases and use all other available rocdl.mfma intrinsics.
Commit: 4813cc5dbaf0b630a5d893ee6f352ac45e3faa40
https://github.com/llvm/llvm-project/commit/4813cc5dbaf0b630a5d893ee6f352ac45e3faa40
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-06 (Mon, 06 Jul 2020)
Changed paths:
A mlir/examples/prototypes/xdlops.mlir
Log Message:
-----------
Add a toy prototype using miopen.mfma.
Commit: 8b865d6a177069b980a56f9809eaad9a9796faec
https://github.com/llvm/llvm-project/commit/8b865d6a177069b980a56f9809eaad9a9796faec
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-07-07 (Tue, 07 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
fixed tuning
Commit: a0bdb16ddddadb579aabcf27080cbf5a3ef7be29
https://github.com/llvm/llvm-project/commit/a0bdb16ddddadb579aabcf27080cbf5a3ef7be29
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-07 (Tue, 07 Jul 2020)
Changed paths:
M mlir/examples/prototypes/xdlops.mlir
M mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
Log Message:
-----------
Experimental commit to support known 9 combinations for fp32 data type.
Commit: f91fbf893e4f7ba45f46ea79d3f9cdd80058d27f
https://github.com/llvm/llvm-project/commit/f91fbf893e4f7ba45f46ea79d3f9cdd80058d27f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-07 (Tue, 07 Jul 2020)
Changed paths:
M mlir/examples/prototypes/xdlops.mlir
M mlir/include/mlir/Dialect/GPU/GPUOps.td
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
M mlir/test/Dialect/GPU/ops.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Remove cbsz, abid, blgp fields from miopen.mfma and gpu.mfma operators.
Commit: 26b61471c9edce5dce767936c10747613867fc64
https://github.com/llvm/llvm-project/commit/26b61471c9edce5dce767936c10747613867fc64
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-07 (Tue, 07 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/GPU/GPUOps.td
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
Log Message:
-----------
Add printer and verifier for gpu.mfma.
Commit: aab131ff488316b63d299f698b4889c110ec49ea
https://github.com/llvm/llvm-project/commit/aab131ff488316b63d299f698b4889c110ec49ea
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/lib/Dialect/Vector/VectorOps.cpp
Log Message:
-----------
XXX relax rule for vector.type_cast.
Need to review this after getting XDLOPS functional.
Commit: f482141aaa51004df00c4be66ec58e95b29b9bc4
https://github.com/llvm/llvm-project/commit/f482141aaa51004df00c4be66ec58e95b29b9bc4
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Revise the syntax of miopen.mfma so it takes memref instead of vector.
Disable lowering test from miopen.mfma to gpu.mfma for now as it needs to be revised.
Commit: 8584eb0493005ebe1b8c8388592b68eb5a16c108
https://github.com/llvm/llvm-project/commit/8584eb0493005ebe1b8c8388592b68eb5a16c108
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
Log Message:
-----------
Revise miopen.mfma to gpu.mfma lowering logic.
Carry out the following tasks in miopen.mfma -> gpu.mfma lowering logic:
- MFMA instruction selection
- MFMA instruction count
- immediates specification
- vector type cast
- vector load / store
Commit: 3207956c9096e0f349e31ba73bbcb9ad954b6ef4
https://github.com/llvm/llvm-project/commit/3207956c9096e0f349e31ba73bbcb9ad954b6ef4
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
M mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
Log Message:
-----------
Revise gpu.mfma to rocdl.mfma lowering logic.
Simplify it as all information are now computed in miopen.mfma -> gpu.mfma.
Commit: 56c7ae0ae4b739ef146a1cffaabdab4efce414ea
https://github.com/llvm/llvm-project/commit/56c7ae0ae4b739ef146a1cffaabdab4efce414ea
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/examples/prototypes/xdlops.mlir
Log Message:
-----------
Revise XDLOPS example code.
Commit: 499c53ccb9bbe848755e9a631a3c5b95a79b0066
https://github.com/llvm/llvm-project/commit/499c53ccb9bbe848755e9a631a3c5b95a79b0066
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
Log Message:
-----------
Fix a bug in fp32 MFMA code selection logic.
Commit: e5c8271cf737ad31cbf7acbecc4e85c04de22326
https://github.com/llvm/llvm-project/commit/e5c8271cf737ad31cbf7acbecc4e85c04de22326
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/examples/prototypes/xdlops.mlir
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Revise syntax of miopen.mfma.
Matrix A and Matrix B types must agree, so just specify one is enough.
Commit: 48253643aa2d4c6c00506d6ba195e961c36447f2
https://github.com/llvm/llvm-project/commit/48253643aa2d4c6c00506d6ba195e961c36447f2
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Extend syntax of miopen.mfma so it takes f16 and bf16.
Commit: cc15043a012cf03a95cdc45888085a0b4e0b41e0
https://github.com/llvm/llvm-project/commit/cc15043a012cf03a95cdc45888085a0b4e0b41e0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
M mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
M mlir/test/Dialect/GPU/ops.mlir
Log Message:
-----------
Revise syntax of gpu.mfma.
Matrix A and Matrix B type must be specified now.
Matrix A and Matrix B types must agree, so just specify one is enough.
Commit: bdd1ef6d5ba4a6696f4d3224d800437f5c9fc83c
https://github.com/llvm/llvm-project/commit/bdd1ef6d5ba4a6696f4d3224d800437f5c9fc83c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/GPU/GPUOps.td
M mlir/test/Dialect/GPU/ops.mlir
Log Message:
-----------
Extend syntax of gpu.mfma so it takes f16 and bf16.
Commit: 8230129426e45a6773884046a0038aa800bbecce
https://github.com/llvm/llvm-project/commit/8230129426e45a6773884046a0038aa800bbecce
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
Log Message:
-----------
Add miopen.mfma -> gpu.mfma lowering logic for FP16.
Commit: 35e2ae47bfe4a60039d4577eb5ff35c60b2b847f
https://github.com/llvm/llvm-project/commit/35e2ae47bfe4a60039d4577eb5ff35c60b2b847f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
M mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
Log Message:
-----------
Add gpu.mfma -> rocdl.mfma lowering logic for FP16.
Commit: 53b94e74cd314a8faaaf7566aac98f73420444fa
https://github.com/llvm/llvm-project/commit/53b94e74cd314a8faaaf7566aac98f73420444fa
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
Log Message:
-----------
Add miopen.mfma -> gpu.mfma lowering logic for BF16.
Commit: eba9286828fe4f014fd8cf9824c2705836858801
https://github.com/llvm/llvm-project/commit/eba9286828fe4f014fd8cf9824c2705836858801
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
M mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
Log Message:
-----------
Add gpu.mfma -> rocdl.mfma lowering logic for BF16.
Commit: 335261680bee174d77bdcdde41eb6f9639c9fc07
https://github.com/llvm/llvm-project/commit/335261680bee174d77bdcdde41eb6f9639c9fc07
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/examples/prototypes/xdlops.mlir
Log Message:
-----------
Revise XDLOPS example code to covert FP16 and BF16 cases.
Commit: dab244197adacf85e6c6d355b790b76a755b77d1
https://github.com/llvm/llvm-project/commit/dab244197adacf85e6c6d355b790b76a755b77d1
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-08 (Wed, 08 Jul 2020)
Changed paths:
M mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
M mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
Log Message:
-----------
Convert from vector<2xbf16> to vector<2xi16> during gpu.mfma -> rocdl.mfma lowering.
This is to cope with types asked by MFMA instrinsics.
Commit: 5b4be26f360f00582bfff2bbe66aad9c1f6c6cca
https://github.com/llvm/llvm-project/commit/5b4be26f360f00582bfff2bbe66aad9c1f6c6cca
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-07-09 (Thu, 09 Jul 2020)
Changed paths:
M mlir/.clang-format
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
fixed comments
Commit: f50f7a620d65af0447b31fd0afe42b94122d15f0
https://github.com/llvm/llvm-project/commit/f50f7a620d65af0447b31fd0afe42b94122d15f0
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-07-09 (Thu, 09 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
Log Message:
-----------
format
Commit: 3798bf955aa885fa7ca44f84164854b5a5291097
https://github.com/llvm/llvm-project/commit/3798bf955aa885fa7ca44f84164854b5a5291097
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-07-09 (Thu, 09 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
Log Message:
-----------
fixed comment
Commit: 976a47dcb4e1dda870c330e3f75fd0e800e2d8cc
https://github.com/llvm/llvm-project/commit/976a47dcb4e1dda870c330e3f75fd0e800e2d8cc
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-13 (Mon, 13 Jul 2020)
Changed paths:
M mlir/.clang-format
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
Log Message:
-----------
Merge pull request #26 from zjing14/tuning_fix
Fixed tuning process
Commit: d2008f23dc8bb65a682cce95959c9023547f614a
https://github.com/llvm/llvm-project/commit/d2008f23dc8bb65a682cce95959c9023547f614a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-17 (Fri, 17 Jul 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Add xdlops attribute in mlir-miopen-driver.
Commit: 20f3e0509fd9e2ba0e7d0ef4284834b7dcf2bba2
https://github.com/llvm/llvm-project/commit/20f3e0509fd9e2ba0e7d0ef4284834b7dcf2bba2
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-17 (Fri, 17 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Add xdlops attribute in miopen.conv2d_* -> miopen.gridwise_gemm lowering.
Commit: 894d86c150e50bceb2994d0fb132fc82cfd157e5
https://github.com/llvm/llvm-project/commit/894d86c150e50bceb2994d0fb132fc82cfd157e5
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-17 (Fri, 17 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Add m_waves and n_waves attributes in miopen.gridwise_gemm -> miopen.blockwise_gemm lowering.
Commit: 989391c7e8d8302a19d4497bf09e031dc61b4a1b
https://github.com/llvm/llvm-project/commit/989391c7e8d8302a19d4497bf09e031dc61b4a1b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-17 (Fri, 17 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Remove non-XDLOPS attributes from XDLOPS lowering pipeline.
Commit: f3bc5e7171a8f410faaee79159015882ccee924a
https://github.com/llvm/llvm-project/commit/f3bc5e7171a8f410faaee79159015882ccee924a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-20 (Mon, 20 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Add xdlops attributes in miopen.gridwise_gemm -> miopen.blockwise_gemm lowering.
Commit: e3c108b1afedf112b34d4a9bf6d1513e755b021a
https://github.com/llvm/llvm-project/commit/e3c108b1afedf112b34d4a9bf6d1513e755b021a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-20 (Mon, 20 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
WIP adding XDLOPS-aware lowering logic for miopen.blockwise_gemm.
Commit: 402d072678a1941d3506eee2e40a06bc7279613b
https://github.com/llvm/llvm-project/commit/402d072678a1941d3506eee2e40a06bc7279613b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-20 (Mon, 20 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Introduce miopen.xdlops_gemm op.
Commit: 34e26faa7b7a7785d3525972e9bde7d6bca3d29f
https://github.com/llvm/llvm-project/commit/34e26faa7b7a7785d3525972e9bde7d6bca3d29f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-20 (Mon, 20 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
WIP adding XDLOPS-aware lowering logic for miopen.blockwise_gemm to miopen.xdlops_gemm.
TBD: coord_transforms attributes are not properly set after lowering miopen.subview.
Commit: 94fa063e680bdfbb585d54d0e4c7ffe84fc1ff63
https://github.com/llvm/llvm-project/commit/94fa063e680bdfbb585d54d0e4c7ffe84fc1ff63
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-24 (Fri, 24 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Refactor blockwise_gemm -> xdlops_gemm lowering logic per xdlops_gemm_repeat_opt branch in MIOpen.
Commit: d4409ce6c9864d6be2352c64c2b916500add6194
https://github.com/llvm/llvm-project/commit/d4409ce6c9864d6be2352c64c2b916500add6194
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-24 (Fri, 24 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Revise syntax of xdlops_gemm, remove offset for Matrix C.
In MIOpen xdlops_gemm_repeat_opt, offset within Matrix C are handled completely
within xdlops_gemm implementation.
Commit: 128d84ab80e8e54baa38a9faee0faf4d8e982b6b
https://github.com/llvm/llvm-project/commit/128d84ab80e8e54baa38a9faee0faf4d8e982b6b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-24 (Fri, 24 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
XXX WIP enhance miopen.subview lowering logic to cope with the requirement of miopen.xdlops_gemm.
miopen.subview for miopen.threadwise_copy we only need the 1st level of affine map.
miopen.subview for miopen.xdlops_gemm, we need to go to the deepest level of affine map.
Commit: 551a0e4f1f97c141131492a164e177ad4a157ad5
https://github.com/llvm/llvm-project/commit/551a0e4f1f97c141131492a164e177ad4a157ad5
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-07-27 (Mon, 27 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
M mlir/tools/mlir-miopen-driver/CMakeLists.txt
A mlir/tools/mlir-miopen-driver/MlirParse.h
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
A mlir/tools/mlir-miopen-driver/mlir-miopen-lib.cpp
Log Message:
-----------
Adding libMLIRMIOpen.a
* Fixed compilation issues exposed in LowerMIOpenOps
* Extracted common file composing logic into MlirParse.h
* Created libMLIRMIOpenThin for standalone static library
* Created libMLIRMIOpen for the target library with LLVM dependencies
Commit: a369fa8201e216bfdb0032edb6aeb332418281cc
https://github.com/llvm/llvm-project/commit/a369fa8201e216bfdb0032edb6aeb332418281cc
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-28 (Tue, 28 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Skeleton of miopen.xdlops_gemm lowering.
Commit: 25481840795d43d60443194ce8ec14369da8d9b0
https://github.com/llvm/llvm-project/commit/25481840795d43d60443194ce8ec14369da8d9b0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-28 (Tue, 28 Jul 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
XXX. Populate hard-coded tuning parameter for XDLOPS setting.
Commit: c4d0f4c94d0a9c9db2b88cb41509c21247d99ad4
https://github.com/llvm/llvm-project/commit/c4d0f4c94d0a9c9db2b88cb41509c21247d99ad4
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-28 (Tue, 28 Jul 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Populate XDLOPS-suitable initial test vector.
Commit: 7c44e01a4770afa315666591aaf40fe053052d9e
https://github.com/llvm/llvm-project/commit/7c44e01a4770afa315666591aaf40fe053052d9e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-28 (Tue, 28 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Continue fill in logic of miopen.xdlops_gemm lowering.
Commit: 3d476a02a41dbe4ce3f5e9d9a991b1b1b7c2c03a
https://github.com/llvm/llvm-project/commit/3d476a02a41dbe4ce3f5e9d9a991b1b1b7c2c03a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-28 (Tue, 28 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
Log Message:
-----------
Amend how Matrix C shall be allocated on XDLOPS path.
Commit: 3816da8dad9ad487c1ad9eabdce6f75383e7860a
https://github.com/llvm/llvm-project/commit/3816da8dad9ad487c1ad9eabdce6f75383e7860a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-28 (Tue, 28 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Continue fill in logic of miopen.xdlops_gemm lowering.
Commit: 9d05e4b86be0f118fd272f861ee2c01a0c1b05e0
https://github.com/llvm/llvm-project/commit/9d05e4b86be0f118fd272f861ee2c01a0c1b05e0
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-28 (Tue, 28 Jul 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
Log Message:
-----------
Allow miopen.mfma to support non-zero address space.
Commit: a647a5a15cc03941460a712298eb0c89e9070f3f
https://github.com/llvm/llvm-project/commit/a647a5a15cc03941460a712298eb0c89e9070f3f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-29 (Wed, 29 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Continue fill in logic of miopen.xdlops_gemm lowering.
Commit: 5f9cc8e6154cc80443936a2f77df8c08f56f7144
https://github.com/llvm/llvm-project/commit/5f9cc8e6154cc80443936a2f77df8c08f56f7144
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-07-29 (Wed, 29 Jul 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/CMakeLists.txt
Log Message:
-----------
Adding fat library and install logic
Commit: cffa1d557e94d19e4e7e9b25828caaefc01237d2
https://github.com/llvm/llvm-project/commit/cffa1d557e94d19e4e7e9b25828caaefc01237d2
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-07-30 (Thu, 30 Jul 2020)
Changed paths:
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Target/CppOutput/ConvertToMIOpenCPP.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/tools/mlir-miopen-driver/CMakeLists.txt
M mlir/tools/mlir-miopen-driver/mlir-miopen-lib.cpp
Log Message:
-----------
Address issues exposed from MIOpen API call
* Get rid of resultStr and update the translate interface
* Make use of OwningModuleRef to guarantee the destruction of module
* Added a genTxt field in the handle that allocates the returned string
Commit: 07e1385f0988294f9a1d654d3fad7022cca0eb77
https://github.com/llvm/llvm-project/commit/07e1385f0988294f9a1d654d3fad7022cca0eb77
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-07-30 (Thu, 30 Jul 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Add XDLOPS matrix C write out skeleton logic.
Commit: 7ac574077ea88094d86eba89176046cb9f7bd3f2
https://github.com/llvm/llvm-project/commit/7ac574077ea88094d86eba89176046cb9f7bd3f2
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-08-02 (Sun, 02 Aug 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
add gemm m/n level tuning
Commit: a644ccedcf1561eca4ad3ff31afb5fa8764c3157
https://github.com/llvm/llvm-project/commit/a644ccedcf1561eca4ad3ff31afb5fa8764c3157
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-08-02 (Sun, 02 Aug 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
clang format
Commit: a3755503dc54ee4bc25f46936d3851091c94bc85
https://github.com/llvm/llvm-project/commit/a3755503dc54ee4bc25f46936d3851091c94bc85
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-02 (Sun, 02 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Extract Xdlops code selection logic from xdlops_gemm lowering logic.
the same logic would be used elsewhere (ex: gridwise_gemm lowering logic).
Commit: ff742a19b58570c93cb1993df36512a208273b3c
https://github.com/llvm/llvm-project/commit/ff742a19b58570c93cb1993df36512a208273b3c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-03 (Mon, 03 Aug 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
Merge pull request #30 from zjing14/add_m_n_cluster_tuning
Added m/n cluster tuning
Commit: 96fd8c445aea21c249047696016bc61609844ba7
https://github.com/llvm/llvm-project/commit/96fd8c445aea21c249047696016bc61609844ba7
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-08-04 (Tue, 04 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
add vector read dim
Commit: 560b5a02d8725e0abd2023f33f7e76de9c46440b
https://github.com/llvm/llvm-project/commit/560b5a02d8725e0abd2023f33f7e76de9c46440b
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-08-04 (Tue, 04 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
clean
Commit: 5afd30adbc4c97d66559e9529cd0ceecc6b52593
https://github.com/llvm/llvm-project/commit/5afd30adbc4c97d66559e9529cd0ceecc6b52593
Author: Jing Zhang <jizhan at amd.com>
Date: 2020-08-04 (Tue, 04 Aug 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
format
Commit: 9bff02461220ffc6b7496796f287d5408a25c8b7
https://github.com/llvm/llvm-project/commit/9bff02461220ffc6b7496796f287d5408a25c8b7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-05 (Wed, 05 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
WIP implementing matrix C write out logic for gridwise_gemm XDLOPS path.
Commit: 88cf7b90ee1ffe11747a5aeca84c54352d61e21e
https://github.com/llvm/llvm-project/commit/88cf7b90ee1ffe11747a5aeca84c54352d61e21e
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-08-05 (Wed, 05 Aug 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/CMakeLists.txt
Log Message:
-----------
Add conditional to build only in static lib scenario
Commit: af84c4874f0d5ce57d44fb8a85e002296a25da31
https://github.com/llvm/llvm-project/commit/af84c4874f0d5ce57d44fb8a85e002296a25da31
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-05 (Wed, 05 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Decouple mMyThreadOffsetA/B with c_thread_mtx_index_row/col.
This paves the way for refactoring XDLOPS lowering logic for blockwise_gemm.
Commit: 3388282762424e66ea34e7761cc43616244a23ea
https://github.com/llvm/llvm-project/commit/3388282762424e66ea34e7761cc43616244a23ea
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-05 (Wed, 05 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
Merge pull request #31 from zjing14/vector_read_dim_tuning
Add Vector read dim tuning
Commit: af1d9360207ed57a65553d322dca4fdf6286262e
https://github.com/llvm/llvm-project/commit/af1d9360207ed57a65553d322dca4fdf6286262e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-05 (Wed, 05 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Rename variable. register C on VGPR could be 1D in XDLOPS path.
Commit: 3c797822674288bcedf398ac6f3a6f0594d4e2f3
https://github.com/llvm/llvm-project/commit/3c797822674288bcedf398ac6f3a6f0594d4e2f3
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-05 (Wed, 05 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Update Matrix C addressing logic.
Original C++ logic uses <M0, 1, M2, 1> for Matrix C on VGPR. And only covers
1/NumBlks part of it.
Updated using <NumBlks, M0, M2> to address it.
Commit: 41550a864a2785391cd9037f6060e2a2e5de6980
https://github.com/llvm/llvm-project/commit/41550a864a2785391cd9037f6060e2a2e5de6980
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-05 (Wed, 05 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
rename ConstantIndexOp -> ConstantOp.
Commit: 1ade15859f79117a668c9df39fc514ca001fb755
https://github.com/llvm/llvm-project/commit/1ade15859f79117a668c9df39fc514ca001fb755
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-06 (Thu, 06 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Revise the layout of matrix C on VGPR.
<NumBlks, M0, M2> -> <1, NumBlks, M0, M2>
Commit: e100ab76d4ea01b491a6208741baec2fc9ccedd7
https://github.com/llvm/llvm-project/commit/e100ab76d4ea01b491a6208741baec2fc9ccedd7
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-08-07 (Fri, 07 Aug 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/CMakeLists.txt
M mlir/test/CMakeLists.txt
A mlir/test/mlir-miopen-lib/populate_bwd.mlir
A mlir/test/mlir-miopen-lib/populate_bww.mlir
A mlir/test/mlir-miopen-lib/populate_fw.mlir
M mlir/tools/mlir-miopen-driver/CMakeLists.txt
A mlir/tools/mlir-miopen-driver/mlir-miopen-lib-test.cpp
M mlir/tools/mlir-miopen-driver/mlir-miopen-lib.cpp
A mlir/tools/mlir-miopen-driver/mlir-miopen-lib.hpp
Log Message:
-----------
Add new unit test target mlir-miopen-lib-test
Commit: ccad4d916b1ad79912d38a5e675cc08a5d369ac8
https://github.com/llvm/llvm-project/commit/ccad4d916b1ad79912d38a5e675cc08a5d369ac8
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-07 (Fri, 07 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Conversion/MIOpenToGPU/CMakeLists.txt
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
M mlir/lib/Target/CppOutput/ConvertToMIOpenCPP.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/test/CMakeLists.txt
A mlir/test/mlir-miopen-lib/populate_bwd.mlir
A mlir/test/mlir-miopen-lib/populate_bww.mlir
A mlir/test/mlir-miopen-lib/populate_fw.mlir
M mlir/tools/mlir-miopen-driver/CMakeLists.txt
A mlir/tools/mlir-miopen-driver/MlirParse.h
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
A mlir/tools/mlir-miopen-driver/mlir-miopen-lib-test.cpp
A mlir/tools/mlir-miopen-driver/mlir-miopen-lib.cpp
A mlir/tools/mlir-miopen-driver/mlir-miopen-lib.hpp
Log Message:
-----------
Merge pull request #29 from jerryyin/miopen-dialect-libMLIRMIOpen
Adding libMLIRMIOpen.a for cpp generation API
Commit: 1931e979d45e0b8d06c2a2639cf397e6d386255f
https://github.com/llvm/llvm-project/commit/1931e979d45e0b8d06c2a2639cf397e6d386255f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-07 (Fri, 07 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Change the syntax of miopen.mfma to take an offset to matrix C.
Commit: 15f41d683166ab55e0f955a72e1843b3418f7cc7
https://github.com/llvm/llvm-project/commit/15f41d683166ab55e0f955a72e1843b3418f7cc7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-07 (Fri, 07 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
A mlir/include/mlir/Dialect/MIOpen/XdlopsCodeSelection.h
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
Log Message:
-----------
Improve xdlops_gemm lowering logic take offset in matrix C into consideration.
Commit: fca12c1bb1afa2e1ffd147c188097c3ebfdd1fd9
https://github.com/llvm/llvm-project/commit/fca12c1bb1afa2e1ffd147c188097c3ebfdd1fd9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-10 (Mon, 10 Aug 2020)
Changed paths:
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
Log Message:
-----------
XXX improve unit test.
Commit: d3590322d5e170bd56789878d3a2bc15c032a05c
https://github.com/llvm/llvm-project/commit/d3590322d5e170bd56789878d3a2bc15c032a05c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-12 (Wed, 12 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Fix matrix C output logic.
Commit: 911b2dac0d0ecca2eeb9c1cbbf53ec646512e26a
https://github.com/llvm/llvm-project/commit/911b2dac0d0ecca2eeb9c1cbbf53ec646512e26a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-13 (Thu, 13 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Fix compile error and LDS offset computation error.
Commit: ed112006ff7471844002c7245520acb4d87c1e0d
https://github.com/llvm/llvm-project/commit/ed112006ff7471844002c7245520acb4d87c1e0d
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-13 (Thu, 13 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
Log Message:
-----------
Revise XDLOPS code selection logic.
Instruction stream is based on MPerXdlops / NPerXdlops, not MPerWave / NPerWave.
Commit: b578450f6da4963e4a5e04a9d71c653831513b94
https://github.com/llvm/llvm-project/commit/b578450f6da4963e4a5e04a9d71c653831513b94
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-13 (Thu, 13 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
XXX commit to get XDLOPS path work end-to-end.
By pass emission MFMAOp for now. Instead, use:
- load the value
- increase by 1 (A: 1, B: 1, A*B = 1)
- store the value
This would allow the first test vector to complete.
Still need to look into why XDLOPS instruction streams are messed up by the optimizer.
Commit: 551832ddebf107543cd31c6ea629cfa1bf89ea98
https://github.com/llvm/llvm-project/commit/551832ddebf107543cd31c6ea629cfa1bf89ea98
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-08-14 (Fri, 14 Aug 2020)
Changed paths:
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/test/Dialect/MIOpen/translate_bwd_data_kcyx_nchw_nkhw.mlir
M mlir/test/Dialect/MIOpen/translate_forward_kcyx_nchw_nkhw.mlir
M mlir/test/mlir-miopen-lib/populate_bwd.mlir
M mlir/test/mlir-miopen-lib/populate_bww.mlir
M mlir/test/mlir-miopen-lib/populate_fw.mlir
Log Message:
-----------
Updating MLIR generated filename and kernel name
Commit: 205cba9df4025e01b2a380b68882411a0482bff3
https://github.com/llvm/llvm-project/commit/205cba9df4025e01b2a380b68882411a0482bff3
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-17 (Mon, 17 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Conversion/MIOpenToGPU/mfma.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Introduce miopen.mfma_v2 op.
Commit: cb55fb7d4fbf910f110150fdd7483989f594656e
https://github.com/llvm/llvm-project/commit/cb55fb7d4fbf910f110150fdd7483989f594656e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-17 (Mon, 17 Aug 2020)
Changed paths:
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
A mlir/test/Conversion/MIOpenToGPU/mfma_v2.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Implement miopen.mfma_v2 to gpu.mfma lowering logic and test.
Commit: 430d9ba20d4aaf5bb669b1792c572156088d96d6
https://github.com/llvm/llvm-project/commit/430d9ba20d4aaf5bb669b1792c572156088d96d6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-17 (Mon, 17 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Introduce miopen.xdlops_gemm_v2 op.
Commit: 072652f149eaaf48dd8e00ec87b0d76cfd22f03a
https://github.com/llvm/llvm-project/commit/072652f149eaaf48dd8e00ec87b0d76cfd22f03a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-18 (Tue, 18 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
A mlir/test/Dialect/MIOpen/lowering_xdlops_gemm_v2.mlir
Log Message:
-----------
Initial commit of miopen.xdlops_gemm_v2 lowering logic and test.
Commit: 3b8d0972ba7a8336ab91d95becda12ae2904b9c9
https://github.com/llvm/llvm-project/commit/3b8d0972ba7a8336ab91d95becda12ae2904b9c9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-18 (Tue, 18 Aug 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Add xdlopsV2 attribute to mlir-miopen-driver.
Commit: 7aa0579b9df93a37b82d9d5ac6528814a64f8cad
https://github.com/llvm/llvm-project/commit/7aa0579b9df93a37b82d9d5ac6528814a64f8cad
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-18 (Tue, 18 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Permeate xdlopsV2 attribute to the lowering process.
Commit: a50ed14fecac1bb5153a0e3dc694fd02e0c0c383
https://github.com/llvm/llvm-project/commit/a50ed14fecac1bb5153a0e3dc694fd02e0c0c383
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-18 (Tue, 18 Aug 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/XdlopsCodeSelection.h
Log Message:
-----------
Add XdlopsCodeEmission logic.
Commit: e582e75cceba159b73abe11581c45ebc3a8d59a1
https://github.com/llvm/llvm-project/commit/e582e75cceba159b73abe11581c45ebc3a8d59a1
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-08-21 (Fri, 21 Aug 2020)
Changed paths:
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/test/Dialect/MIOpen/translate_bwd_data_kcyx_nchw_nkhw.mlir
M mlir/test/Dialect/MIOpen/translate_forward_kcyx_nchw_nkhw.mlir
M mlir/test/mlir-miopen-lib/populate_bwd.mlir
M mlir/test/mlir-miopen-lib/populate_bww.mlir
M mlir/test/mlir-miopen-lib/populate_fw.mlir
Log Message:
-----------
Merge pull request #32 from jerryyin/miopen-dialect
Updating MLIR generated filename and kernel name
Commit: 4915716e819cf20f267d4b4e04308e7b37b95a45
https://github.com/llvm/llvm-project/commit/4915716e819cf20f267d4b4e04308e7b37b95a45
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-08 (Tue, 08 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/lowering_xdlops_gemm_v2.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Revise syntax of miopen.xdlops_gemm_v2 to allow variadic input and output.
Commit: ded872dcd0b64c1213a5270c698ff5a2495dd125
https://github.com/llvm/llvm-project/commit/ded872dcd0b64c1213a5270c698ff5a2495dd125
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-08 (Tue, 08 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/XdlopsCodeSelection.h
Log Message:
-----------
Revise XdlopsCodeSelection logic to keep with latest MIOpen implementation.
Commit: 48afaf6a35b22b7a2659cc62aa50b74250b4ba47
https://github.com/llvm/llvm-project/commit/48afaf6a35b22b7a2659cc62aa50b74250b4ba47
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-08 (Tue, 08 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/XdlopsCodeSelection.h
Log Message:
-----------
Merge XdlopsCodeEmission logic into XdlopsCodeSelection.
Commit: 1c25c7d498b4d1cfad90cad786c7eb88b908ba26
https://github.com/llvm/llvm-project/commit/1c25c7d498b4d1cfad90cad786c7eb88b908ba26
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-08 (Tue, 08 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/test/Dialect/MIOpen/lowering_xdlops_gemm_v2.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Revised implementation of miopen.xdlops_gemm_v2.
Commit: 0be11cf8e54792c27c4d8ecc73bca7055ab47440
https://github.com/llvm/llvm-project/commit/0be11cf8e54792c27c4d8ecc73bca7055ab47440
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-08 (Tue, 08 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Add miopen.blockwise_gemm_v2 op and test.
Commit: 78252424c776d7a205a490a8c6b12b6c7e65f5d1
https://github.com/llvm/llvm-project/commit/78252424c776d7a205a490a8c6b12b6c7e65f5d1
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-08 (Tue, 08 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Initial commit of skeleton of blockwise_gemm_v2 lowering logic.
Commit: b7bd26c61be358a00d0279995631f55fdc2f1804
https://github.com/llvm/llvm-project/commit/b7bd26c61be358a00d0279995631f55fdc2f1804
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-08 (Tue, 08 Sep 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
Log Message:
-----------
Improve miopen.blockwise_gemm_v2 & miopen.xdlops_gemm_v2 print logic.
Commit: c3f2b964e08ba6491d4a989000ee153fc560cefd
https://github.com/llvm/llvm-project/commit/c3f2b964e08ba6491d4a989000ee153fc560cefd
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-08 (Tue, 08 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Implementation of miopen.blockwise_gemm_v2 lowering logic.
- Considers MRepeats = NRepeats = 1 case.
- Considers MRepeats = 1, NRepeats = 2 case.
- Considers MRepeats = 2, NRepeats = 1 case.
- Add attributes.
Commit: 017ab453624873888a1128c64c0ea5ab0f23159c
https://github.com/llvm/llvm-project/commit/017ab453624873888a1128c64c0ea5ab0f23159c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-08 (Tue, 08 Sep 2020)
Changed paths:
A mlir/test/Dialect/MIOpen/lowering_blockwise_gemm_v2.mlir
Log Message:
-----------
Add test for lowering logic for miopen.blockwise_gemm_v2.
Commit: a8851d0ca5df692ad71dcab76127d2654a904c4f
https://github.com/llvm/llvm-project/commit/a8851d0ca5df692ad71dcab76127d2654a904c4f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-09 (Wed, 09 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_level.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Remove unused miopen.gridwise_gemm_ex operation and its tests.
Commit: 2aa073645c9151f71553131f000ec08157808b5b
https://github.com/llvm/llvm-project/commit/2aa073645c9151f71553131f000ec08157808b5b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-09 (Wed, 09 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops.mlir
Log Message:
-----------
Initial skeleton of miopen.gridwise_gemm_v2.
This operation would contain the logic of XDLOPS-specific gridwise_gemm.
Commit: 78c6e00599412b651f8a851a57301b244e1b002c
https://github.com/llvm/llvm-project/commit/78c6e00599412b651f8a851a57301b244e1b002c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-09 (Wed, 09 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Add miopen.gridwise_gemm_v2 into miope.conv2d* lowering pipeline.
Emit miopen.gridwise_gemm_v2 when xdlopsV2 attribute is set true.
Commit: 6af0ebe71f05d4eb2da82aa19418ba252762131c
https://github.com/llvm/llvm-project/commit/6af0ebe71f05d4eb2da82aa19418ba252762131c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-09 (Wed, 09 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
Let -miopen-affix-params pass support miopen.gridwise_gemm_v2 op.
Commit: ce7508fb7e9c2e88fbd3708e9021d21c1c704fe1
https://github.com/llvm/llvm-project/commit/ce7508fb7e9c2e88fbd3708e9021d21c1c704fe1
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-09 (Wed, 09 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Skeleton implementation of miopen.gridwise_gemm_v2 lowering logic.
Simply copy all lowering logic of miopen.gridwise_gemm for now.
Next steps:
- allocate C matrices as vectors instead of memrefs.
- invoke miopen.blockwise_gemm_v2 with vectors.
- change miopen.threadwise_copy_v2 so it accept vectors.
- invoke miopen.threadwise_copy_v2.
- end-to-end test on fp32.
Further next steps:
- Remove obsolete XDLOPS path from miopen.gridwise_gemm logic.
- Revise lower-level parts (miopen.mfma, miopen.xdlops_gemm_v2, miopen.blockwise_gemm_v2) for fp16 and bf16.
- Revise -miopen-affine-transform pass to take fp16 and bf16 into consideration.
- Revise -miopen-affix-params pass to take fp16 and bf16 into consideration.
- Revise miopen.conv2d* rewriting logic to take fp16 and bf16 into consideration.
- Revise miopen.gridwise_gemm_v2 lowering logic to adopt MIOpen GridwiseBatchGemmXdlops_gkmkpack_gknkpack_gmn_v2.
Commit: a33ed627289a57dd165986792f0334e36f5e1fbf
https://github.com/llvm/llvm-project/commit/a33ed627289a57dd165986792f0334e36f5e1fbf
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-09 (Wed, 09 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Define miopen.threadwise_copy_v2 for matrix C write out.
Commit: 2a9689cf8eb962072060d24854e75cf501d54ada
https://github.com/llvm/llvm-project/commit/2a9689cf8eb962072060d24854e75cf501d54ada
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-09 (Wed, 09 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
A mlir/test/Dialect/MIOpen/lowering_threadwise_copy_v2.mlir
Log Message:
-----------
Skeleton implementation of miopen.threadwise_copy_v2 lowering logic and test.
Commit: e6a4ce854d432f67caa7489432da5e254b856027
https://github.com/llvm/llvm-project/commit/e6a4ce854d432f67caa7489432da5e254b856027
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-10 (Thu, 10 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/test/Dialect/MIOpen/lowering_threadwise_copy_v2.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Implementation of miopen.threadwise_copy_v2 lowering logic.
Added more test cases.
Commit: 32a887fd829c40d5563ea1df31d4c16ea9cd6934
https://github.com/llvm/llvm-project/commit/32a887fd829c40d5563ea1df31d4c16ea9cd6934
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-11 (Fri, 11 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Work-in-progress miopen.gridwise_gemm_v2 lowering logic.
- Add XDLOPS code selection.
- Allocate vectors on C and 0 initialize them.
- Invoke miopen.blockwise_gemm_v2 operations in LDS double buffer loop.
- Set attributes for miopen.blockwise_gemm_v2 via information inside miopen.gridwise_gemm_v2.
- Implement matrix C write out logic using miopen.threadwise_copy_v2.
- Set attributes for miopen.threadwise_gemm_v2 via information inside miopen.gridwise_gemm_v2.
TBD:
- Set attributes for miopen.xdlops_gemm_v2 via information inside miopen.blockwise_gemm_v2.
Commit: d6bd952ac3b4f5d820f944a0b9b7b48e3e8cf55c
https://github.com/llvm/llvm-project/commit/d6bd952ac3b4f5d820f944a0b9b7b48e3e8cf55c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-11 (Fri, 11 Sep 2020)
Changed paths:
A mlir/examples/prototypes/test_xdlops_v2.mlir
Log Message:
-----------
Add an example prototype of requirements for gridwise_gemm_v2.
Commit: 12e4d12f138688a2b47ffd8618880eabff9bb761
https://github.com/llvm/llvm-project/commit/12e4d12f138688a2b47ffd8618880eabff9bb761
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-11 (Fri, 11 Sep 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Move miopen.xdlops_gemm_v2 lowering to -miopen-lowering-step5.
Commit: f8b009001816e0420eabbffcf60c1de06236cddc
https://github.com/llvm/llvm-project/commit/f8b009001816e0420eabbffcf60c1de06236cddc
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-11 (Fri, 11 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/test/Dialect/MIOpen/lowering_xdlops_gemm_v2.mlir
Log Message:
-----------
XXX WIP tuning miopen.xdlops_gemm_v2 lowering logic.
Commit: e15b6ff6515d2001b9b3d1029bf0123cff666057
https://github.com/llvm/llvm-project/commit/e15b6ff6515d2001b9b3d1029bf0123cff666057
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-11 (Fri, 11 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/XdlopsCodeSelection.h
Log Message:
-----------
Fix XDLOPS code selection logic for fp32.
Commit: 5627a62520d92baf384ade727ca0c4bbd7c3a8f2
https://github.com/llvm/llvm-project/commit/5627a62520d92baf384ade727ca0c4bbd7c3a8f2
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-09-14 (Mon, 14 Sep 2020)
Changed paths:
M mlir/tools/mlir-rocm-runner/CMakeLists.txt
M mlir/utils/jenkins/Jenkinsfile.downstream
Log Message:
-----------
Bump Jenkins to ROCm3.7 and amdhip64
Commit: c3c7c0fb4e3facb19625ddc3a14fc3b1bbb0035a
https://github.com/llvm/llvm-project/commit/c3c7c0fb4e3facb19625ddc3a14fc3b1bbb0035a
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-09-14 (Mon, 14 Sep 2020)
Changed paths:
M mlir/utils/jenkins/Jenkinsfile.downstream
Log Message:
-----------
Removing unecessary checkout stage
Commit: bc3a304a9d92b66b37dc4f5371005fdfb2627e8b
https://github.com/llvm/llvm-project/commit/bc3a304a9d92b66b37dc4f5371005fdfb2627e8b
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-09-14 (Mon, 14 Sep 2020)
Changed paths:
M mlir/utils/jenkins/Jenkinsfile.downstream
Log Message:
-----------
Cleanup workspace after build
Commit: be8d9dd1eee19004da6e085bdd06c352b8900c88
https://github.com/llvm/llvm-project/commit/be8d9dd1eee19004da6e085bdd06c352b8900c88
Author: Zhuoran Yin <zhuoryin at amd.com>
Date: 2020-09-14 (Mon, 14 Sep 2020)
Changed paths:
M mlir/tools/mlir-rocm-runner/CMakeLists.txt
M mlir/utils/jenkins/Jenkinsfile.downstream
Log Message:
-----------
Merge pull request #3 from ROCmSoftwarePlatform/jenkins-rocm3.7
Bump Jenkins to ROCm3.7 and amdhip64
Commit: 24670360de5b0fbb5f9a849bd378bfa931fb86b7
https://github.com/llvm/llvm-project/commit/24670360de5b0fbb5f9a849bd378bfa931fb86b7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-17 (Thu, 17 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
Log Message:
-----------
WIP: attempt to reduce stack usage by moving allocation inside xdlops_gemm_v2 to gridwise_gemm_v2.
Commit: d9bd73307ee99dd7e95ea0665d50a1a21b5f967f
https://github.com/llvm/llvm-project/commit/d9bd73307ee99dd7e95ea0665d50a1a21b5f967f
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-18 (Fri, 18 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Fix vector C offset computation logic.
Commit: 0f7de27c481132f5f464a497c6dece87330c2ab6
https://github.com/llvm/llvm-project/commit/0f7de27c481132f5f464a497c6dece87330c2ab6
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-23 (Wed, 23 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
Log Message:
-----------
Got matrix C write out logic correct.
Every place of output tensor would now be written.
Commit: 1875399ff01ff6b6a03b0821076dc0e186a41ec4
https://github.com/llvm/llvm-project/commit/1875399ff01ff6b6a03b0821076dc0e186a41ec4
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-23 (Wed, 23 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Reduce LDS size by half. Remove LDS double buffer loop.
For XDLOPS path we don't need to do LDS double buffering.
The codes are disabled for now and will be enabled in later commits.
Also refactor codes.
Commit: 9633d598f628cd8335032bdef7a11ba3e0b96d21
https://github.com/llvm/llvm-project/commit/9633d598f628cd8335032bdef7a11ba3e0b96d21
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-23 (Wed, 23 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Re-enable XDLOPS GEMM for just 1 iteration at the loop tail.
Commit: 898a41a90db0fdd91abb0875d0749dbbe9198655
https://github.com/llvm/llvm-project/commit/898a41a90db0fdd91abb0875d0749dbbe9198655
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-23 (Wed, 23 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Remove obsolete comments.
Commit: 2860ad870b96109c6cea060f8047d5948b47027b
https://github.com/llvm/llvm-project/commit/2860ad870b96109c6cea060f8047d5948b47027b
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/lib/Conversion/GPUCommon/CMakeLists.txt
M mlir/lib/Conversion/GPUCommon/ConvertKernelFuncToBlob.cpp
Log Message:
-----------
XXX use -O0 when generating ISA for GPU kernels.
Commit: 3d319f08d3578e9bd560a3c0f0c107e960ac6a46
https://github.com/llvm/llvm-project/commit/3d319f08d3578e9bd560a3c0f0c107e960ac6a46
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Enable MFMA hot loop.
LDS blockwise_copy is not enabled yet.
Commit: 0bb3f1bf0092e1438b7aafaed28e25fc5d9e7a4c
https://github.com/llvm/llvm-project/commit/0bb3f1bf0092e1438b7aafaed28e25fc5d9e7a4c
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/lib/Conversion/GPUCommon/CMakeLists.txt
M mlir/lib/Conversion/GPUCommon/ConvertKernelFuncToBlob.cpp
M mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp
Log Message:
-----------
Change amdgpu-flat-work-group-size to 1,256 to enable -O3 optimizations.
Commit: a81d05c42edf4749e2ff52d8debe77a67095d3f2
https://github.com/llvm/llvm-project/commit/a81d05c42edf4749e2ff52d8debe77a67095d3f2
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
Log Message:
-----------
Compute gridSize properly instead of hard-coding.
Commit: 8d371ddb4de87758aa0b8f5ad096f21a33ed0689
https://github.com/llvm/llvm-project/commit/8d371ddb4de87758aa0b8f5ad096f21a33ed0689
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Remove obsolete comments.
Commit: 50ea5fb450b3e1bec02c63fc7cf7981da2cf6bb7
https://github.com/llvm/llvm-project/commit/50ea5fb450b3e1bec02c63fc7cf7981da2cf6bb7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Remove miopen.xdlops_gemm.
miopen.xdlops_gemm_v2 is getting closer to fully functional.
Remove miopen.xdlops_gemm.
Commit: d1b7802b74035259b71ce3a092a7e88cd7e75a0a
https://github.com/llvm/llvm-project/commit/d1b7802b74035259b71ce3a092a7e88cd7e75a0a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
R mlir/test/Conversion/MIOpenToGPU/mfma.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Remove miopen.mfma.
miopen.mfma_v2 is getting closer to fully functional.
Remove miopen.mfma and its relevant logic and tests.
Commit: 8d9171a3a672ba72080dcd03a22d5a58512ea600
https://github.com/llvm/llvm-project/commit/8d9171a3a672ba72080dcd03a22d5a58512ea600
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/test/Dialect/MIOpen/ops_2.mlir
Log Message:
-----------
Fix unit tests for miopen.blockwise_gemm_v2 and miopen.xdlops_gemm_v2.
Commit: 721f187b0a3deb34d14a80035498cf1a97c9d311
https://github.com/llvm/llvm-project/commit/721f187b0a3deb34d14a80035498cf1a97c9d311
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Remove obsolete xdlops (V1) lowering pipeline.
Commit: 360bf511e7e07c116b26bf158d26e6ff8aa8d755
https://github.com/llvm/llvm-project/commit/360bf511e7e07c116b26bf158d26e6ff8aa8d755
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Enable blockwise_copy and loading from LDS in MFMA lowering.
Commit: 7d9b86e0979ae1ef08b67a40251efba04c5e0c83
https://github.com/llvm/llvm-project/commit/7d9b86e0979ae1ef08b67a40251efba04c5e0c83
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-24 (Thu, 24 Sep 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Enable additional logic in miopen.xdlops_gemm_v2 lowering logic.
Commit: f17a260c711f60dd79dac2e474eafa32ccae193d
https://github.com/llvm/llvm-project/commit/f17a260c711f60dd79dac2e474eafa32ccae193d
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-09-25 (Fri, 25 Sep 2020)
Changed paths:
A mlir/examples/prototypes/step2_try_use_vector_for_blockwise_copy_coords.mlir
Log Message:
-----------
Add one prototype which try to use vector types for blockwise_copy coordinates.
Commit: 81ab4863e57e5673ad9307c38a7292b5f6ba4ae7
https://github.com/llvm/llvm-project/commit/81ab4863e57e5673ad9307c38a7292b5f6ba4ae7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-01 (Thu, 01 Oct 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Allow fp16 / bf16 types in mlir-miopen-driver.
Commit: ea31e3b39892e1d58e2a8c8368212bb2cbcc8ba7
https://github.com/llvm/llvm-project/commit/ea31e3b39892e1d58e2a8c8368212bb2cbcc8ba7
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-01 (Thu, 01 Oct 2020)
Changed paths:
A mlir/examples/prototypes/step2_try_use_vector_for_blockwise_copy_coords.mlir
Log Message:
-----------
Merge branch 'miopen-dialect-xdlops-take2' of github.com:whchung/llvm-project into miopen-dialect-xdlops-take2
Commit: a4eace7e88bab9aba0bfa262a06f7980085b7b42
https://github.com/llvm/llvm-project/commit/a4eace7e88bab9aba0bfa262a06f7980085b7b42
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-01 (Thu, 01 Oct 2020)
Changed paths:
M mlir/tools/mlir-rocm-runner/CMakeLists.txt
Log Message:
-----------
Migrate to hip-clang runtime library.
Commit: b0ab6307e50317e6795f719b22e0c988c84b1d3e
https://github.com/llvm/llvm-project/commit/b0ab6307e50317e6795f719b22e0c988c84b1d3e
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-01 (Thu, 01 Oct 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
Log Message:
-----------
Amend the dialect so fp16/bf16 types are supported at all levels.
Logic to properly address fp16/bf16 are NOT implemented yet.
Commit: 9b91945d51986743cf652fcae119609038479829
https://github.com/llvm/llvm-project/commit/9b91945d51986743cf652fcae119609038479829
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-05 (Mon, 05 Oct 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/XdlopsCodeSelection.h
Log Message:
-----------
Add logic in XDLOPS code selection to specify type for argA / argB.
Commit: 2b2b11fd64e647e86d2308d1178cb99f4bd05366
https://github.com/llvm/llvm-project/commit/2b2b11fd64e647e86d2308d1178cb99f4bd05366
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-05 (Mon, 05 Oct 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Getting fp16 and bf16 support into the lowering pipeline.
Commit: fca59b401f0955b6613e303deb029d30211766eb
https://github.com/llvm/llvm-project/commit/fca59b401f0955b6613e303deb029d30211766eb
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-06 (Tue, 06 Oct 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Modify mlir-miopen-driver to populate host logic for fp16 and bf16.
Commit: ef67e7da006f29b8533342b06900845cce4d6a6a
https://github.com/llvm/llvm-project/commit/ef67e7da006f29b8533342b06900845cce4d6a6a
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-06 (Tue, 06 Oct 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
M mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp
Log Message:
-----------
Let HIP runtime wrappers support fp16 and bf16 types.
Commit: 5d1de5f175b8398eee7db824af6246fd0eb5f8eb
https://github.com/llvm/llvm-project/commit/5d1de5f175b8398eee7db824af6246fd0eb5f8eb
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-06 (Tue, 06 Oct 2020)
Changed paths:
M mlir/include/mlir/ExecutionEngine/RunnerUtils.h
M mlir/lib/ExecutionEngine/RunnerUtils.cpp
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Work-in-progress print fp16/bf16 memref logic.
Commit: 8984f06707d1d4f34f8fd3cd5b14770f38f003a4
https://github.com/llvm/llvm-project/commit/8984f06707d1d4f34f8fd3cd5b14770f38f003a4
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-07 (Wed, 07 Oct 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Teach host population logic to support backward data and backward weight.
Commit: 173e00dcf3b7b22acca1935b07682777a2b32204
https://github.com/llvm/llvm-project/commit/173e00dcf3b7b22acca1935b07682777a2b32204
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-07 (Wed, 07 Oct 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
Log Message:
-----------
Fix tensor coordinate transformation for backward weight.
Commit: 7449419b94f2d00c7c1378f57b3d465b8f2c80b8
https://github.com/llvm/llvm-project/commit/7449419b94f2d00c7c1378f57b3d465b8f2c80b8
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-08 (Thu, 08 Oct 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/include/mlir/Target/MIOpenCPP.h
M mlir/lib/Conversion/MIOpenToGPU/CMakeLists.txt
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
M mlir/lib/Target/CppOutput/ConvertToMIOpenCPP.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/test/CMakeLists.txt
M mlir/test/Dialect/MIOpen/translate_bwd_data_kcyx_nchw_nkhw.mlir
M mlir/test/Dialect/MIOpen/translate_forward_kcyx_nchw_nkhw.mlir
A mlir/test/mlir-miopen-lib/populate_bwd.mlir
A mlir/test/mlir-miopen-lib/populate_bww.mlir
A mlir/test/mlir-miopen-lib/populate_fw.mlir
M mlir/tools/mlir-miopen-driver/CMakeLists.txt
A mlir/tools/mlir-miopen-driver/MlirParse.h
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
A mlir/tools/mlir-miopen-driver/mlir-miopen-lib-test.cpp
A mlir/tools/mlir-miopen-driver/mlir-miopen-lib.cpp
A mlir/tools/mlir-miopen-driver/mlir-miopen-lib.hpp
M mlir/utils/jenkins/Jenkinsfile.downstream
Log Message:
-----------
Merge branch 'miopen-dialect' into miopen-dialect-xdlops-take2
Commit: 8c0fffe97a88144efa55a37046eb93e70155ba16
https://github.com/llvm/llvm-project/commit/8c0fffe97a88144efa55a37046eb93e70155ba16
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-12 (Mon, 12 Oct 2020)
Changed paths:
A mlir/examples/prototypes/step2_try_use_vector_for_blockwise_copy_coords.mlir
A mlir/examples/prototypes/test_xdlops_v2.mlir
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
A mlir/include/mlir/Dialect/MIOpen/XdlopsCodeSelection.h
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/include/mlir/ExecutionEngine/RunnerUtils.h
M mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp
M mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
M mlir/lib/ExecutionEngine/RunnerUtils.cpp
M mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp
R mlir/test/Conversion/MIOpenToGPU/mfma.mlir
A mlir/test/Conversion/MIOpenToGPU/mfma_v2.mlir
A mlir/test/Dialect/MIOpen/lowering_blockwise_gemm_v2.mlir
M mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_level.mlir
A mlir/test/Dialect/MIOpen/lowering_threadwise_copy_v2.mlir
A mlir/test/Dialect/MIOpen/lowering_xdlops_gemm_v2.mlir
M mlir/test/Dialect/MIOpen/ops.mlir
M mlir/test/Dialect/MIOpen/ops_2.mlir
M mlir/tools/mlir-miopen-driver/MlirParse.h
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
M mlir/tools/mlir-miopen-driver/mlir-miopen-lib.cpp
M mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp
Log Message:
-----------
Merge pull request #5 from ROCmSoftwarePlatform/miopen-dialect-xdlops-take2
XDLOPS support logic
Commit: 00bbf5fd0730ff5d73d9425cb70f2a7208afb2d1
https://github.com/llvm/llvm-project/commit/00bbf5fd0730ff5d73d9425cb70f2a7208afb2d1
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-10-12 (Mon, 12 Oct 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
M mlir/test/Dialect/MIOpen/lowering_filter_tensor_ckyx_cnhw_knhw.mlir
M mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_position_cyxk_chwn_khwn.mlir
M mlir/test/Dialect/MIOpen/lowering_input_tensor_cyxk_cnhw_hnhw.mlir
M mlir/test/Dialect/MIOpen/lowering_memref_kcyx_nchw_nkhw.mlir
M mlir/test/Dialect/MIOpen/lowering_output_tensor_kyxc_nhwc_nhwk.mlir
M mlir/test/Dialect/MIOpen/lowering_top_level.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_bwd.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_fw.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
M mlir/test/mlir-miopen-driver/populate.mlir
M mlir/test/mlir-miopen-driver/populate_host.mlir
M mlir/test/mlir-miopen-lib/populate_bwd.mlir
M mlir/test/mlir-miopen-lib/populate_bww.mlir
M mlir/test/mlir-miopen-lib/populate_fw.mlir
M mlir/tools/mlir-miopen-driver/CMakeLists.txt
M mlir/tools/mlir-miopen-driver/MlirParse.h
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
M mlir/tools/mlir-miopen-driver/mlir-miopen-lib.cpp
Log Message:
-----------
Populate arch and cu to convolution context
Commit: c3289a4d2507919ca509d2e034b55b8cd1637493
https://github.com/llvm/llvm-project/commit/c3289a4d2507919ca509d2e034b55b8cd1637493
Author: jerryyin <zhuoryin at amd.com>
Date: 2020-10-12 (Mon, 12 Oct 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/CMakeLists.txt
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Address review feedback - get rid of runtime dependency
Commit: 93bf0ebe9c196b4a723d8d1edfdb319545b78cc9
https://github.com/llvm/llvm-project/commit/93bf0ebe9c196b4a723d8d1edfdb319545b78cc9
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-13 (Tue, 13 Oct 2020)
Changed paths:
M mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
M mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
M mlir/test/Dialect/MIOpen/lowering_filter_tensor_ckyx_cnhw_knhw.mlir
M mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_position_cyxk_chwn_khwn.mlir
M mlir/test/Dialect/MIOpen/lowering_input_tensor_cyxk_cnhw_hnhw.mlir
M mlir/test/Dialect/MIOpen/lowering_memref_kcyx_nchw_nkhw.mlir
M mlir/test/Dialect/MIOpen/lowering_output_tensor_kyxc_nhwc_nhwk.mlir
M mlir/test/Dialect/MIOpen/lowering_top_level.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_bwd.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_fw.mlir
M mlir/test/Dialect/MIOpen/translate_cflags_wrw.mlir
M mlir/test/mlir-miopen-driver/populate.mlir
M mlir/test/mlir-miopen-driver/populate_host.mlir
M mlir/test/mlir-miopen-lib/populate_bwd.mlir
M mlir/test/mlir-miopen-lib/populate_bww.mlir
M mlir/test/mlir-miopen-lib/populate_fw.mlir
M mlir/tools/mlir-miopen-driver/CMakeLists.txt
M mlir/tools/mlir-miopen-driver/MlirParse.h
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
M mlir/tools/mlir-miopen-driver/mlir-miopen-lib.cpp
Log Message:
-----------
Merge pull request #7 from ROCmSoftwarePlatform/populate_arch_cu
Populate arch and cu to convolution context
Commit: e7a4b323209dc19fd50cd10ae8dd37af29f49183
https://github.com/llvm/llvm-project/commit/e7a4b323209dc19fd50cd10ae8dd37af29f49183
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-19 (Mon, 19 Oct 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
A mlir/utils/jenkins/Jenkinsfile.perf
Log Message:
-----------
performance benchmarking Jenkinsfile step1 (#12)
* Add a flag to supress result tensor print out logic.
* Initial perf benchmarking Jenkinsfile.
Commit: 3237ce8afc6af771bca36870cd0c48ff5b8febd6
https://github.com/llvm/llvm-project/commit/3237ce8afc6af771bca36870cd0c48ff5b8febd6
Author: Zhuoran Yin <zhuoryin at amd.com>
Date: 2020-10-19 (Mon, 19 Oct 2020)
Changed paths:
A mlir/include/mlir/Dialect/MIOpen/Tuning/GridwiseGemmParams.h
A mlir/include/mlir/Dialect/MIOpen/Tuning/Serializable.h
A mlir/include/mlir/Dialect/MIOpen/Tuning/SqliteDb.h
R mlir/include/mlir/Dialect/MIOpen/gridwise_gemm_params.h
M mlir/lib/Dialect/MIOpen/CMakeLists.txt
M mlir/lib/Dialect/MIOpen/Transforms/AffixTuningParameters.cpp
M mlir/lib/Dialect/MIOpen/Transforms/CMakeLists.txt
A mlir/lib/Dialect/MIOpen/Tuning/CMakeLists.txt
A mlir/lib/Dialect/MIOpen/Tuning/GridwiseGemmParams.cpp
A mlir/lib/Dialect/MIOpen/Tuning/SqliteDb.cpp
A mlir/lib/Dialect/MIOpen/Tuning/cmake/FindSQLite3.cmake
M mlir/lib/Target/CppOutput/CMakeLists.txt
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
A mlir/test/Dialect/MIOpen/Tuning/db_connection.mlir
A mlir/test/Dialect/MIOpen/Tuning/db_load_fail.mlir
A mlir/test/Dialect/MIOpen/Tuning/db_load_success.mlir
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
M mlir/tools/mlir-rocm-runner/CMakeLists.txt
M mlir/utils/jenkins/Jenkinsfile.downstream
Log Message:
-----------
Perfdb client to query tuning parameters (#8)
* Implementing MLIR perfdb client
* Install miopen for perfdb
* Addressing review feedback part1
* Moved cmake change from mlir/CmakeLists to mlir/lib/Dialect/MIOpen to
make it scoped
* Renamed source and header to be consistent with rest of MLIR
* Added Tuning directory for header
* Add sqlite db tests
* Install and pull in sqlite3 without external dependencies
Commit: 48bf37f48fb709b70e9f62856d7d62a1bc9b69ab
https://github.com/llvm/llvm-project/commit/48bf37f48fb709b70e9f62856d7d62a1bc9b69ab
Author: yiqian1 <68618497+yiqian1 at users.noreply.github.com>
Date: 2020-10-20 (Tue, 20 Oct 2020)
Changed paths:
M mlir/utils/jenkins/Jenkinsfile.downstream
A mlir/utils/jenkins/static-checks/clang-format.ignore
A mlir/utils/jenkins/static-checks/clang-tidy.ignore
A mlir/utils/jenkins/static-checks/premerge-checks.py
Log Message:
-----------
Add clang-format and clang-tidy checks into Jenkins Pipeline (#9)
* Add clang-format and clang-tidy checks into Jenkins Pipeline
* Move premerge-check scripts to a directory and update the jenkinsfile
* Remove a docker-run argument which caused a conflict when installing python3-pip. And fix a typo
* Invokde premerge-checks.py with python3.6
* Fix missused single quotes
* Remove exit code and add comments
* Fix error
Co-authored-by: Yi Qian <[yi.qian at amd.com]>
Commit: 176aefc1fc0fb359c7c9a7b4674c8bc92d1db2bb
https://github.com/llvm/llvm-project/commit/176aefc1fc0fb359c7c9a7b4674c8bc92d1db2bb
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-20 (Tue, 20 Oct 2020)
Changed paths:
M mlir/lib/Dialect/MIOpen/Tuning/CMakeLists.txt
Log Message:
-----------
Cosmetic change in SQLite config process. (#13)
Commit: 54da861e0e6268b21f5c05ab8e4da054f5a23198
https://github.com/llvm/llvm-project/commit/54da861e0e6268b21f5c05ab8e4da054f5a23198
Author: Zhuoran Yin <zhuoryin at amd.com>
Date: 2020-10-20 (Tue, 20 Oct 2020)
Changed paths:
M mlir/utils/jenkins/Jenkinsfile.downstream
R mlir/utils/jenkins/static-checks/clang-format.ignore
R mlir/utils/jenkins/static-checks/clang-tidy.ignore
R mlir/utils/jenkins/static-checks/premerge-checks.py
Log Message:
-----------
Revert "Add clang-format and clang-tidy checks into Jenkins Pipeline (#9)" (#15)
This reverts commit 48bf37f48fb709b70e9f62856d7d62a1bc9b69ab.
Commit: ae0b7825c738b95cf517b55678fa9d5d1146e694
https://github.com/llvm/llvm-project/commit/ae0b7825c738b95cf517b55678fa9d5d1146e694
Author: Zhuoran Yin <zhuoryin at amd.com>
Date: 2020-10-20 (Tue, 20 Oct 2020)
Changed paths:
A mlir/include/mlir/Dialect/MIOpen/Tuning/ConvContext.h
M mlir/include/mlir/Dialect/MIOpen/Tuning/GridwiseGemmParams.h
M mlir/lib/Dialect/MIOpen/Tuning/GridwiseGemmParams.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
M mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4_gen_xdlops.cpp
Log Message:
-----------
Refactor to bring XDLOps Tuning up to date (#14)
Commit: 3911342c8d3e22a1e62b7fdce463faae0e01ebc3
https://github.com/llvm/llvm-project/commit/3911342c8d3e22a1e62b7fdce463faae0e01ebc3
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-20 (Tue, 20 Oct 2020)
Changed paths:
M mlir/tools/mlir-miopen-driver/mlir-miopen-driver.cpp
Log Message:
-----------
Populate default arch and num_cu for gfx908. (#16)
Commit: efb4c85ae698de9beae0d1a191d023c6805763f4
https://github.com/llvm/llvm-project/commit/efb4c85ae698de9beae0d1a191d023c6805763f4
Author: yiqian1 <68618497+yiqian1 at users.noreply.github.com>
Date: 2020-10-21 (Wed, 21 Oct 2020)
Changed paths:
M mlir/utils/jenkins/Jenkinsfile.downstream
A mlir/utils/jenkins/static-checks/clang-format.ignore
A mlir/utils/jenkins/static-checks/clang-tidy.ignore
A mlir/utils/jenkins/static-checks/premerge-checks.py
Log Message:
-----------
Test static checks (#18)
* Add clang-format and clang-tidy checks into Jenkins Pipeline
* Add clang-format and clang-tidy checks into Jenkins Pipeline
Commit: 72ab085bad56c7d86eb1cd6ee0eb2a85d7ee5c67
https://github.com/llvm/llvm-project/commit/72ab085bad56c7d86eb1cd6ee0eb2a85d7ee5c67
Author: Wen-Heng (Jack) Chung <whchung at gmail.com>
Date: 2020-10-22 (Thu, 22 Oct 2020)
Changed paths:
M mlir/utils/jenkins/Jenkinsfile.perf
Log Message:
-----------
Improve performance benchmarking jenkins script.
Ensure SQLite unit tests pass.
Compare: https://github.com/llvm/llvm-project/compare/125143fddbef%5E...72ab085bad56
More information about the All-commits
mailing list