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

Impl Kthvalue operation #3204

Merged
merged 31 commits into from
Oct 1, 2024
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
31 commits
Select commit Hold shift + click to select a range
e154cca
rebase skeleton code with develop
BuiChiTrung Aug 20, 2024
6e9898c
add fwd kernel with driver
BuiChiTrung Jun 26, 2024
b27cb59
check improvement over rocm
BuiChiTrung Jun 27, 2024
0f0e59f
validate tensor length in prob description
BuiChiTrung Jun 28, 2024
7853ddb
add keepDim param
BuiChiTrung Jul 1, 2024
c290a2e
backup code
BuiChiTrung Jul 2, 2024
e4ffcec
add constraint to accept only last dim
BuiChiTrung Jul 2, 2024
f30dddd
add utils function to tensor.cpp
BuiChiTrung Jul 3, 2024
3b7ac60
minor changes
BuiChiTrung Jul 3, 2024
e6707dd
apply the kernel to non-cont tensor
BuiChiTrung Jul 4, 2024
9865951
remove redundant include preprocessors
BuiChiTrung Jul 5, 2024
58693f3
backup
BuiChiTrung Jul 5, 2024
c3d229a
resolve comments
BuiChiTrung Jul 5, 2024
a552ba2
resolve cmts
BuiChiTrung Jul 5, 2024
8ddbb85
fix cmts
BuiChiTrung Jul 5, 2024
ab4ee73
resolve cmts
BuiChiTrung Jul 8, 2024
81af7f3
update include header
BuiChiTrung Jul 9, 2024
aec61d2
remove outdate functions
BuiChiTrung Aug 20, 2024
c485922
resolve comments and fix window build fail
BuiChiTrung Aug 21, 2024
bc382a2
clang-format
BuiChiTrung Aug 21, 2024
df0fbc9
fix case isnan compile err in fp16 tests
BuiChiTrung Aug 22, 2024
ef36a8f
merge upstream/develop
BuiChiTrung Aug 22, 2024
732cead
update test name config
BuiChiTrung Aug 23, 2024
be322dd
update tensor view util to pass cppcheck
BuiChiTrung Aug 23, 2024
dcc4d76
try comment out unit-test to see if the pipeline pass
BuiChiTrung Aug 23, 2024
452e1c1
fix gtest format
BuiChiTrung Aug 23, 2024
664afa4
check hip tidy
BuiChiTrung Aug 26, 2024
eedd52b
update kernel comment
BuiChiTrung Aug 31, 2024
8baccb8
rollback src/include/miopen/solver/implicitgemm_ck_util.hpp
BuiChiTrung Sep 25, 2024
5683600
Merge remote-tracking branch 'upstream/develop' into impl_kth_value
BuiChiTrung Sep 25, 2024
64d6d11
Merge branch 'develop' into impl_kth_value
junliume Sep 27, 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 @@ -37,4 +37,5 @@ The MIOpen API library is structured as follows:
* :doc:`ReduceCalculation <../doxygen/html/group__ReduceCalculation>` (experimental)
* :doc:`RotaryPositionalEmbeddings <../doxygen/html/group__RotaryPositionalEmbeddings>` (experimental)
* :doc:`ReLU <../doxygen/html/group___re_l_u>` (experimental)
* :doc:`Kthvalue <../doxygen/html/group__kthvalue>` (experimental)
* :doc:`GLU <../doxygen/html/group__glu>` (experimental)
1 change: 1 addition & 0 deletions driver/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ add_executable(MIOpenDriver
dm_getitem.cpp
dm_glu.cpp
dm_groupnorm.cpp
dm_kthvalue.cpp
dm_layernorm.cpp
dm_lrn.cpp
dm_pool.cpp
Expand Down
41 changes: 41 additions & 0 deletions driver/dm_kthvalue.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,41 @@
/*******************************************************************************
*
* 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 "kthvalue_driver.hpp"

static Driver* makeDriver(const std::string& base_arg)
{
if(base_arg == "kthvalue")
return new KthvalueDriver<float>();
else if(base_arg == "kthvaluefp16")
return new KthvalueDriver<float16>();
else if(base_arg == "kthvaluebfp16")
return new KthvalueDriver<bfloat16>();
return nullptr;
}

REGISTER_DRIVER_MAKER(makeDriver);
5 changes: 3 additions & 2 deletions driver/driver.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -176,7 +176,7 @@ inline void PadBufferSize(size_t& sz, int datatype_sz)
"t5layernorm[bfp16|fp16], adam[fp16], ampadam, reduceextreme[bfp16|fp16], "
"adamw[fp16], ampadamw, transformersadamw[fp16], transformersampadamw, "
"getitem[bfp16|fp16], reducecalculation[bfp16|fp16], rope[bfp16|fp16], "
"prelu[bfp16|fp16], glu[bfp16|fp16]\n");
"prelu[bfp16|fp16], kthvalue[bfp16|fp16], glu[bfp16|fp16]\n");
exit(0); // NOLINT (concurrency-mt-unsafe)
}

Expand Down Expand Up @@ -209,7 +209,8 @@ inline std::string ParseBaseArg(int argc, char* argv[])
arg != "getitemfp16" && arg != "getitembfp16" && arg != "reducecalculation" &&
arg != "reducecalculationfp16" && arg != "reducecalculationbfp16" && arg != "rope" &&
arg != "ropefp16" && arg != "ropebfp16" && arg != "prelu" && arg != "prelufp16" &&
arg != "prelubfp16" && arg != "glu" && arg != "glufp16" && arg != "glubfp16" &&
arg != "prelubfp16" && arg != "kthvalue" && arg != "kthvaluefp16" &&
arg != "kthvaluebfp16" && arg != "glu" && arg != "glufp16" && arg != "glubfp16" &&
arg != "--version")
{
printf("FAILED: Invalid Base Input Argument\n");
Expand Down
Loading