Skip to content
New issue

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

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

Already on GitHub? Sign in to your account

Implement Fold and Unfold #3167

Open
wants to merge 52 commits into
base: develop
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from 46 commits
Commits
Show all changes
52 commits
Select commit Hold shift + click to select a range
113620a
UnfoldFwd4d driver, test and api
DuongQLee Jul 2, 2024
6902fdf
githook format
DuongQLee Jul 4, 2024
0f15ed5
githook format
DuongQLee Jul 4, 2024
9b49b9d
unfold backward driver and gtest
DuongQLee Jul 4, 2024
83fba56
githook format
DuongQLee Jul 4, 2024
101794f
Add foldfwd, foldbwd, problem_description verification, gtest and driver
DuongQLee Jul 10, 2024
9286ce7
githook format
DuongQLee Jul 10, 2024
e59ce36
update doc and miopen.h description
DuongQLee Jul 12, 2024
c114934
Update driver help text
DuongQLee Jul 15, 2024
bd5db59
Change IN_OUT_TYPE to FLOAT
DuongQLee Jul 15, 2024
4bb5855
add __restrict__ to tensor pointer
DuongQLee Jul 15, 2024
918a267
replace include "" with <>
DuongQLee Jul 17, 2024
4f51b6e
change all int -> int32_t, remove duplicate lines in solver
DuongQLee Jul 17, 2024
e726fc1
githook format
DuongQLee Jul 17, 2024
8ee2861
remove useless if else in problem description
DuongQLee Jul 18, 2024
f3dea16
add more tensor_layout_t constructor and update kernel to use get_ten…
DuongQLee Jul 18, 2024
4a83296
githook format
DuongQLee Jul 18, 2024
299117b
remove {}
DuongQLee Jul 18, 2024
e76095e
update code as comments
DuongQLee Jul 22, 2024
526f772
githook format
DuongQLee Jul 22, 2024
2a3d2b0
cpu_fold -> cpu_unfold
DuongQLee Jul 22, 2024
27e26c3
update code as comments
DuongQLee Jul 22, 2024
366e350
githook format
DuongQLee Jul 22, 2024
8f14828
Merge branch 'develop' into mv_510_fold
DuongQLee Jul 24, 2024
2b3bd1f
githook format
DuongQLee Jul 24, 2024
5c506b4
Update gtest code
DuongQLee Jul 29, 2024
4e97559
Merge remote-tracking branch 'origin' into mv_510_fold
DuongQLee Jul 29, 2024
e2631d8
githook format
DuongQLee Jul 29, 2024
99beed8
Merge remote-tracking branch 'origin' into mv_510_fold, update code f…
DuongQLee Jul 31, 2024
7bb583b
git hook format
DuongQLee Jul 31, 2024
a6256e7
add MIOPEN_INTERNALS_EXPORT
DuongQLee Jul 31, 2024
d359c1f
githook format
DuongQLee Jul 31, 2024
80424c6
Merge remote-tracking branch 'origin' into mv_510_fold
DuongQLee Aug 2, 2024
1652dc4
resolve conflict
DuongQLee Aug 2, 2024
1f6e4a2
githook format
DuongQLee Aug 2, 2024
2f9bce7
fix git merge dup
DuongQLee Aug 2, 2024
fa6f15a
githook format
DuongQLee Aug 2, 2024
45ed5c1
update tensor_view and kernel code
DuongQLee Aug 5, 2024
ba2020b
githook format
DuongQLee Aug 5, 2024
87edcbd
remove duplicate miopen ops and update doc
DuongQLee Aug 5, 2024
db7b9a8
update spacing
DuongQLee Aug 5, 2024
857db5c
empty commit
DuongQLee Aug 6, 2024
0da1cc6
update gtest syntax
DuongQLee Aug 6, 2024
879c5c7
githook format
DuongQLee Aug 6, 2024
9e66c78
Merge branch 'develop' into mv_510_fold
DuongQLee Aug 12, 2024
7d07012
githook format
DuongQLee Aug 12, 2024
85c1ee0
add not contiguous test cases for fold and unfold
DuongQLee Aug 13, 2024
66e5dcb
githook format
DuongQLee Aug 13, 2024
918091d
remove /*context*/ for solver
DuongQLee Aug 13, 2024
901d7b3
remove gen_one
DuongQLee Aug 13, 2024
8956248
githook format
DuongQLee Aug 13, 2024
c4e1d12
Merge branch 'develop' into impl_fold_unfold
long10024070 Aug 19, 2024
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions docs/reference/index.rst
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,7 @@ The MIOpen API library is structured as follows:
* :doc:`Cat <../doxygen/html/group__cat>` (experimental)
* :doc:`SGD <../doxygen/html/group___s_g_d>` (experimental)
* :doc:`ReduceExtreme <../doxygen/html/group__ReduceExtreme>` (experimental)
* :doc:`Fold <./group___f_o_l_d>` (experimental)
* :doc:`Getitem <../doxygen/html/group__getitem>` (experimental)
* :doc:`ReduceCalculation <../doxygen/html/group__ReduceCalculation>` (experimental)
* :doc:`RotaryPositionalEmbeddings <../doxygen/html/group__RotaryPositionalEmbeddings>` (experimental)
2 changes: 2 additions & 0 deletions driver/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -57,6 +57,8 @@ add_executable(MIOpenDriver
dm_t5layernorm.cpp
dm_tensorop.cpp
dm_transformers_adam_w.cpp
dm_fold.cpp
dm_unfold.cpp
main.cpp
registry_driver_maker.cpp
rocrand_wrapper.cpp)
Expand Down
39 changes: 39 additions & 0 deletions driver/dm_fold.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,39 @@
/*******************************************************************************
*
* MIT License
*
* Copyright (c) 2024 Advanced Micro Devices, Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*
*******************************************************************************/
#include "registry_driver_maker.hpp"
#include "fold_driver.hpp"
static Driver* makeDriver(const std::string& base_arg)
{
if(base_arg == "fold")
return new FoldDriver<float, float>();
if(base_arg == "foldfp16")
return new FoldDriver<float16, float>();
if(base_arg == "foldbfp16")
return new FoldDriver<bfloat16, float>();
return nullptr;
}

REGISTER_DRIVER_MAKER(makeDriver);
39 changes: 39 additions & 0 deletions driver/dm_unfold.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,39 @@
/*******************************************************************************
*
* MIT License
*
* Copyright (c) 2024 Advanced Micro Devices, Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*
*******************************************************************************/
#include "registry_driver_maker.hpp"
#include "unfold_driver.hpp"
static Driver* makeDriver(const std::string& base_arg)
{
if(base_arg == "unfold")
return new UnfoldDriver<float, float>();
if(base_arg == "unfoldfp16")
return new UnfoldDriver<float16, float>();
if(base_arg == "unfoldbfp16")
return new UnfoldDriver<bfloat16, float>();
return nullptr;
}

REGISTER_DRIVER_MAKER(makeDriver);
20 changes: 12 additions & 8 deletions driver/driver.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -169,13 +169,15 @@ inline void PadBufferSize(size_t& sz, int datatype_sz)
[[noreturn]] inline void Usage()
{
printf("Usage: ./driver *base_arg* *other_args*\n");
printf("Supported Base Arguments: conv[fp16|int8|bfp16], pool[fp16], lrn[fp16], "
"activ[fp16], softmax[fp16], bnorm[fp16], rnn[fp16], gemm[fp16], ctc, dropout[fp16], "
"tensorop, reduce[fp16|fp64], layernorm[bfp16|fp16], sum[bfp16|fp16], "
"groupnorm[bfp16|fp16], cat[bfp16|fp16], addlayernorm[bfp16|fp16], "
"t5layernorm[bfp16|fp16], adam[fp16], ampadam, reduceextreme[bfp16|fp16], "
"adamw[fp16], ampadamw, transformersadamw[fp16], transformersampadamw, "
"getitem[bfp16|fp16], reducecalculation[bfp16|fp16], rope[bfp16|fp16]\n");
printf(
"Supported Base Arguments: conv[fp16|int8|bfp16], pool[fp16], lrn[fp16], "
"activ[fp16], softmax[fp16], bnorm[fp16], rnn[fp16], gemm[fp16], ctc, dropout[fp16], "
"tensorop, reduce[fp16|fp64], layernorm[bfp16|fp16], sum[bfp16|fp16], "
"groupnorm[bfp16|fp16], cat[bfp16|fp16], addlayernorm[bfp16|fp16], "
"t5layernorm[bfp16|fp16], adam[fp16], ampadam, reduceextreme[bfp16|fp16], "
"adamw[fp16], ampadamw, transformersadamw[fp16], transformersampadamw, "
"getitem[bfp16|fp16], reducecalculation[bfp16|fp16], rope[bfp16|fp16], unfold[bfp16|fp16], "
"fold[bfp16|fp16]\n");
exit(0); // NOLINT (concurrency-mt-unsafe)
}

Expand Down Expand Up @@ -207,7 +209,9 @@ inline std::string ParseBaseArg(int argc, char* argv[])
arg != "transformersadamwfp16" && arg != "transformersampadamw" && arg != "getitem" &&
arg != "getitemfp16" && arg != "getitembfp16" && arg != "reducecalculation" &&
arg != "reducecalculationfp16" && arg != "reducecalculationbfp16" && arg != "rope" &&
arg != "ropefp16" && arg != "ropebfp16" && arg != "--version")
arg != "ropefp16" && arg != "ropebfp16" && arg != "unfold" && arg != "unfoldfp16" &&
arg != "unfoldbfp16" && arg != "fold" && arg != "foldfp16" && arg != "foldbfp16" &&
arg != "--version")
{
printf("FAILED: Invalid Base Input Argument\n");
Usage();
Expand Down
Loading