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 SmoothL1Loss #3348

Open
wants to merge 91 commits into
base: develop
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from 77 commits
Commits
Show all changes
91 commits
Select commit Hold shift + click to select a range
ca9c098
Adding Smooth L1Loss Forward API - Unreduction kernel, without verify…
long10024070 Apr 10, 2024
61ef667
update license 2023 to 2024
long10024070 Apr 10, 2024
db8b31c
add SmoothL1Loss GTest and debug miopenSmoothL1LossForward api
long10024070 Apr 10, 2024
ec2dabe
using same testcases for all reduction mode in SmoothL1Loss GTest
long10024070 Apr 10, 2024
ac8b114
using float instead of float* for Beta param in SmoothL1Loss api
long10024070 Apr 10, 2024
ea602af
Using 'N,C,D,H,W' for describe SmoothL1Loss testcases' tensors
long10024070 Apr 11, 2024
de8e427
shorten code
long10024070 Apr 11, 2024
1267085
added SmoothL1Loss to MIOpenDriver
long10024070 Apr 11, 2024
2cfd171
remove .gitignore
long10024070 Apr 12, 2024
54876ca
move TEST_P outside smooth_l1loss
long10024070 Apr 12, 2024
f2cb4db
moved new things from SmoothL1Loss to the last positions
long10024070 Apr 12, 2024
6472292
fixed DOXYGEN GROUP from BackendAPI to LossFunction
long10024070 Apr 12, 2024
1e6a439
Merge with commit 2fdf814 from Develop branch, resolved Driver code c…
long10024070 Apr 12, 2024
d91ddfc
Merge branch 'develop-moreh' into impl_SmoothL1Loss
long10024070 Apr 12, 2024
d295e15
change HIP kernel format'
long10024070 Apr 16, 2024
2dfab74
added test for half and bfloat16
long10024070 Apr 16, 2024
4d1618c
applied Clang format
long10024070 Apr 16, 2024
634ffbd
tunned LOCAL_SIZE
long10024070 Apr 16, 2024
40d5a02
added fixed testcases
long10024070 Apr 16, 2024
3e3e372
remove unused lines of code
long10024070 Apr 16, 2024
7460768
Used the same error comparing method with driver in SmoothL1Loss Gtest
long10024070 Apr 16, 2024
51da6bf
remove beta from SmoothL1Loss cache label
long10024070 Apr 16, 2024
281509c
formated code
long10024070 Apr 16, 2024
e80adc0
update documentation
long10024070 Apr 16, 2024
08230ea
changed SmoothL1Loss GTest env requiremnt
long10024070 Apr 16, 2024
8d18be8
corrected ordering
long10024070 Apr 17, 2024
4d64d75
mre data for kernel indentification
long10024070 Apr 17, 2024
423aee5
changed namespace and directory from to
long10024070 Apr 17, 2024
9b985b0
changed namespace and directory from to
long10024070 Apr 17, 2024
644751b
shorten code
long10024070 Apr 22, 2024
8df2dcc
included the operator's name
long10024070 Apr 22, 2024
837b231
removed IsCorrectReduction() check function
long10024070 Apr 22, 2024
a98e01d
remove smooth_l1loss.rst
long10024070 Apr 22, 2024
b5e9448
ported SmoothL1LossUnreducedForward5d kernel
long10024070 Apr 25, 2024
f99eb06
add GTest for SmoothL1LossUnreducedForwardNoncontiguous
long10024070 Apr 26, 2024
48989a3
added SmoothL1LossUnreducedForward5d Driver
long10024070 Apr 26, 2024
7e03c71
debug and reformat SmoothL1Loss driver
long10024070 Apr 27, 2024
47dae9a
debug and reformat SmoothL1Loss driver
long10024070 Apr 27, 2024
3216f39
Changed Driver format and tunning BlockSize
long10024070 Apr 29, 2024
811af8e
Isolated SmoothL1LossUnreduced with reduction cases
long10024070 Apr 29, 2024
e5a78d6
Implemented SmoothL1LossReducedForward5d
long10024070 May 2, 2024
bb2e94a
Removed SameType checker and and defined output type as FLOATTING typ…
long10024070 May 2, 2024
75d8eb9
code shorten
long10024070 May 2, 2024
291de86
debug elapsed time measurement
long10024070 May 2, 2024
0ebd01c
improved SmoothL1LossReduceForward performance
long10024070 May 2, 2024
9beb2a8
removed DTYPE equivalent between Input and Target tensors
long10024070 May 3, 2024
227bac4
make SmoothL1LossReducedForward5d and LossSum kernels to be compiled …
long10024070 May 3, 2024
92394b7
Impl SmoothL1LossReducedBackward5d
long10024070 May 8, 2024
0bc0a2d
Impl SmoothL1LossReducedBackward5d
long10024070 May 8, 2024
d55a89e
changed target tensor array using INPUT_TYPE
long10024070 May 8, 2024
a9eb747
Impl SmoothL1LossUnreducedBackwardContiguous
long10024070 May 8, 2024
3a9e38e
removed non improved over Rocm operation
long10024070 May 8, 2024
e5025d0
updated solver.cpp
long10024070 May 9, 2024
ff61a06
changed MIOPEN_THROW calling position
long10024070 May 9, 2024
fd6fb08
decreased generator range for improving precision and fixing data ove…
long10024070 May 12, 2024
ca65e5c
Fixed workspace allocation bug in Driver
long10024070 May 12, 2024
1092fa9
removed unnecessary strides check
long10024070 May 15, 2024
682e690
Small update for GTest
long10024070 Oct 7, 2024
ee971f9
Merge branch 'impl_SmoothL1Loss' into `develop`
long10024070 Oct 7, 2024
5349def
update code to latest commit (to be continue)
long10024070 Oct 8, 2024
9d53ae8
update smooth_l1loss_driver #include header
long10024070 Oct 21, 2024
a547eb5
update branch
long10024070 Oct 31, 2024
259a371
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Oct 31, 2024
b5a2a9c
sort CMakeLists to lexorder
long10024070 Oct 31, 2024
58f65c6
update for cppcheck
long10024070 Oct 31, 2024
94381f2
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Nov 1, 2024
85a623d
fix Backward kernel bugs
long10024070 Nov 1, 2024
b122ad1
consisten loss reduce sum
long10024070 Nov 1, 2024
2b2d7ac
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Nov 1, 2024
77c828c
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Nov 1, 2024
90a873a
Merge remote-tracking branch 'origin/impl_SmoothL1Loss' into impl_Smo…
long10024070 Nov 1, 2024
6e94c06
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Nov 4, 2024
a14c8ed
remove unused header
long10024070 Nov 4, 2024
0f862c7
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Nov 5, 2024
66cda3d
update doxygen
long10024070 Nov 5, 2024
0049710
Merge remote-tracking branch 'rocm/impl_SmoothL1Loss' into impl_Smoot…
long10024070 Nov 5, 2024
1e0cda1
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Nov 18, 2024
08c3f11
Update test/gtest/smoothl1loss.hpp
long10024070 Nov 19, 2024
f9dfe7c
Update driver/smooth_l1loss_driver.hpp
long10024070 Nov 19, 2024
4332a9f
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Nov 19, 2024
90b8176
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Nov 25, 2024
5a01364
fix for clang format
long10024070 Nov 25, 2024
dd23c9c
add `install_dir/` to .gitignore
long10024070 Nov 25, 2024
6f06768
Merge branch 'impl_SmoothL1Loss' into HEAD
long10024070 Nov 25, 2024
cd1ba52
update
long10024070 Nov 25, 2024
cd7c979
rollback .gitignore
long10024070 Nov 26, 2024
b5c7402
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Dec 1, 2024
109e65d
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Dec 2, 2024
ee0cdcc
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Dec 9, 2024
7392200
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Dec 11, 2024
4087b7d
Merge branch 'develop' into impl_SmoothL1Loss
long10024070 Dec 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 @@ -39,3 +39,4 @@ The MIOpen API library is structured as follows:
* :doc:`ReLU <../doxygen/html/group___re_l_u>` (experimental)
* :doc:`Kthvalue <../doxygen/html/group__kthvalue>` (experimental)
* :doc:`GLU <../doxygen/html/group__glu>` (experimental)
* :doc:`SmoothL1Loss <../doxygen/html/_smooth_l1_loss>` (experimental)
1 change: 1 addition & 0 deletions driver/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -57,6 +57,7 @@ add_executable(MIOpenDriver
dm_reducecalculation.cpp
dm_rnn.cpp
dm_rope.cpp
dm_smooth_l1loss.cpp
dm_softmarginloss.cpp
dm_softmax.cpp
dm_t5layernorm.cpp
Expand Down
41 changes: 41 additions & 0 deletions driver/dm_smooth_l1loss.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 "smooth_l1loss_driver.hpp"

static Driver* makeDriver(const std::string& base_arg)
{
if(base_arg == "smoothl1loss")
return new SmoothL1LossDriver<float, float>();
if(base_arg == "smoothl1lossfp16")
return new SmoothL1LossDriver<float16, float>();
if(base_arg == "smoothl1lossbfp16")
return new SmoothL1LossDriver<bfloat16, float>();
return nullptr;
}

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

Expand Down Expand Up @@ -352,6 +352,7 @@ inline std::string ParseBaseArg(int argc, char* argv[])
arg != "kthvaluebfp16" && arg != "glu" && arg != "glufp16" && arg != "glubfp16" &&
arg != "softmarginloss" && arg != "softmarginlossfp16" && arg != "softmarginlossbfp16" &&
arg != "multimarginloss" && arg != "multimarginlossfp16" && arg != "multimarginlossbfp16" &&
arg != "smoothl1loss" && arg != "smoothl1lossfp16" && arg != "smoothl1lossbfp16" &&
arg != "--version")
{
printf("FAILED: Invalid Base Input Argument\n");
Expand Down
Loading