Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel_ROCM
Commits
8ae8a558
Commit
8ae8a558
authored
Feb 29, 2024
by
aska-0096
Browse files
Merge branch 'develop' of
https://github.com/ROCmSoftwarePlatform/composable_kernel
into navi3_rel
parents
0c0ddeff
a776978c
Changes
66
Show whitespace changes
Inline
Side-by-side
Showing
20 changed files
with
31 additions
and
29 deletions
+31
-29
client_example/01_gemm/gemm.cpp
client_example/01_gemm/gemm.cpp
+2
-2
client_example/02_gemm_add_add_fastgelu/gemm_add_add_fastgelu.cpp
...xample/02_gemm_add_add_fastgelu/gemm_add_add_fastgelu.cpp
+2
-2
client_example/02_gemm_add_add_fastgelu/gemm_add_add_fastgelu_generic.cpp
...2_gemm_add_add_fastgelu/gemm_add_add_fastgelu_generic.cpp
+2
-2
client_example/02_gemm_add_add_fastgelu/gemm_add_fastgelu.cpp
...nt_example/02_gemm_add_add_fastgelu/gemm_add_fastgelu.cpp
+2
-2
client_example/02_gemm_add_add_fastgelu/gemm_add_fastgelu_generic.cpp
...le/02_gemm_add_add_fastgelu/gemm_add_fastgelu_generic.cpp
+2
-2
client_example/02_gemm_add_add_fastgelu/gemm_fastgelu.cpp
client_example/02_gemm_add_add_fastgelu/gemm_fastgelu.cpp
+2
-2
client_example/02_gemm_add_add_fastgelu/gemm_fastgelu_generic.cpp
...xample/02_gemm_add_add_fastgelu/gemm_fastgelu_generic.cpp
+2
-2
client_example/03_gemm_layernorm/gemm_add_add_layernorm_naive.cpp
...xample/03_gemm_layernorm/gemm_add_add_layernorm_naive.cpp
+4
-2
client_example/03_gemm_layernorm/gemm_add_relu_add_layernorm_welford.cpp
...03_gemm_layernorm/gemm_add_relu_add_layernorm_welford.cpp
+2
-2
client_example/04_contraction/contraction_bilinear_fp32.cpp
client_example/04_contraction/contraction_bilinear_fp32.cpp
+1
-1
client_example/04_contraction/contraction_bilinear_fp64.cpp
client_example/04_contraction/contraction_bilinear_fp64.cpp
+1
-1
client_example/04_contraction/contraction_g1m2n3k1_add_xdl_fp16.cpp
...mple/04_contraction/contraction_g1m2n3k1_add_xdl_fp16.cpp
+1
-1
client_example/04_contraction/contraction_scale_fp32.cpp
client_example/04_contraction/contraction_scale_fp32.cpp
+1
-1
client_example/04_contraction/contraction_scale_fp64.cpp
client_example/04_contraction/contraction_scale_fp64.cpp
+1
-1
client_example/05_layernorm/layernorm2d_bwd_data.cpp
client_example/05_layernorm/layernorm2d_bwd_data.cpp
+1
-1
client_example/05_layernorm/layernorm2d_bwd_gamma_beta.cpp
client_example/05_layernorm/layernorm2d_bwd_gamma_beta.cpp
+1
-1
client_example/05_layernorm/layernorm2d_fwd.cpp
client_example/05_layernorm/layernorm2d_fwd.cpp
+1
-1
client_example/05_layernorm/layernorm4d_fwd.cpp
client_example/05_layernorm/layernorm4d_fwd.cpp
+1
-1
client_example/06_softmax/softmax4d.cpp
client_example/06_softmax/softmax4d.cpp
+1
-1
client_example/07_grouped_convnd_fwd/grouped_conv1d_fwd.cpp
client_example/07_grouped_convnd_fwd/grouped_conv1d_fwd.cpp
+1
-1
No files found.
client_example/01_gemm/gemm.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
@@ -83,7 +83,7 @@ int main(int argc, char* argv[])
[](
std
::
size_t
nRow
,
std
::
size_t
nCol
,
std
::
size_t
stride
,
auto
layout
)
{
using
Layout
=
decltype
(
layout
);
if
constexpr
(
std
::
is_same
<
Layout
,
ck
::
tensor_layout
::
gemm
::
RowMajor
>::
value
)
if
constexpr
(
std
::
is_same
<
Layout
,
Row
>::
value
)
{
return
(
nRow
-
1
)
*
stride
+
nCol
;
}
...
...
client_example/02_gemm_add_add_fastgelu/gemm_add_add_fastgelu.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
@@ -92,7 +92,7 @@ int main(int argc, char* argv[])
[](
std
::
size_t
nRow
,
std
::
size_t
nCol
,
std
::
size_t
stride
,
auto
layout
)
{
using
Layout
=
decltype
(
layout
);
if
constexpr
(
std
::
is_same
<
Layout
,
ck
::
tensor_layout
::
gemm
::
RowMajor
>::
value
)
if
constexpr
(
std
::
is_same
<
Layout
,
Row
>::
value
)
{
return
(
nRow
-
1
)
*
stride
+
nCol
;
}
...
...
client_example/02_gemm_add_add_fastgelu/gemm_add_add_fastgelu_generic.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
@@ -93,7 +93,7 @@ int main(int argc, char* argv[])
[](
std
::
size_t
nRow
,
std
::
size_t
nCol
,
std
::
size_t
stride
,
auto
layout
)
{
using
Layout
=
decltype
(
layout
);
if
constexpr
(
std
::
is_same
<
Layout
,
ck
::
tensor_layout
::
gemm
::
RowMajor
>::
value
)
if
constexpr
(
std
::
is_same
<
Layout
,
Row
>::
value
)
{
return
(
nRow
-
1
)
*
stride
+
nCol
;
}
...
...
client_example/02_gemm_add_add_fastgelu/gemm_add_fastgelu.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
@@ -88,7 +88,7 @@ int main(int argc, char* argv[])
[](
std
::
size_t
nRow
,
std
::
size_t
nCol
,
std
::
size_t
stride
,
auto
layout
)
{
using
Layout
=
decltype
(
layout
);
if
constexpr
(
std
::
is_same
<
Layout
,
ck
::
tensor_layout
::
gemm
::
RowMajor
>::
value
)
if
constexpr
(
std
::
is_same
<
Layout
,
Row
>::
value
)
{
return
(
nRow
-
1
)
*
stride
+
nCol
;
}
...
...
client_example/02_gemm_add_add_fastgelu/gemm_add_fastgelu_generic.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
@@ -89,7 +89,7 @@ int main(int argc, char* argv[])
[](
std
::
size_t
nRow
,
std
::
size_t
nCol
,
std
::
size_t
stride
,
auto
layout
)
{
using
Layout
=
decltype
(
layout
);
if
constexpr
(
std
::
is_same
<
Layout
,
ck
::
tensor_layout
::
gemm
::
RowMajor
>::
value
)
if
constexpr
(
std
::
is_same
<
Layout
,
Row
>::
value
)
{
return
(
nRow
-
1
)
*
stride
+
nCol
;
}
...
...
client_example/02_gemm_add_add_fastgelu/gemm_fastgelu.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
@@ -84,7 +84,7 @@ int main(int argc, char* argv[])
[](
std
::
size_t
nRow
,
std
::
size_t
nCol
,
std
::
size_t
stride
,
auto
layout
)
{
using
Layout
=
decltype
(
layout
);
if
constexpr
(
std
::
is_same
<
Layout
,
ck
::
tensor_layout
::
gemm
::
RowMajor
>::
value
)
if
constexpr
(
std
::
is_same
<
Layout
,
Row
>::
value
)
{
return
(
nRow
-
1
)
*
stride
+
nCol
;
}
...
...
client_example/02_gemm_add_add_fastgelu/gemm_fastgelu_generic.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
@@ -85,7 +85,7 @@ int main(int argc, char* argv[])
[](
std
::
size_t
nRow
,
std
::
size_t
nCol
,
std
::
size_t
stride
,
auto
layout
)
{
using
Layout
=
decltype
(
layout
);
if
constexpr
(
std
::
is_same
<
Layout
,
ck
::
tensor_layout
::
gemm
::
RowMajor
>::
value
)
if
constexpr
(
std
::
is_same
<
Layout
,
Row
>::
value
)
{
return
(
nRow
-
1
)
*
stride
+
nCol
;
}
...
...
client_example/03_gemm_layernorm/gemm_add_add_layernorm_naive.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
@@ -17,6 +17,8 @@
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
Row
=
ck
::
tensor_layout
::
gemm
::
RowMajor
;
using
ADataType
=
F16
;
using
BDataType
=
F16
;
using
BiasDataType
=
F32
;
...
...
@@ -191,7 +193,7 @@ int main()
[](
std
::
size_t
nRow
,
std
::
size_t
nCol
,
std
::
size_t
stride
,
auto
layout
)
{
using
Layout
=
decltype
(
layout
);
if
constexpr
(
std
::
is_same
<
Layout
,
ck
::
tensor_layout
::
gemm
::
RowMajor
>::
value
)
if
constexpr
(
std
::
is_same
<
Layout
,
Row
>::
value
)
{
return
(
nRow
-
1
)
*
stride
+
nCol
;
}
...
...
client_example/03_gemm_layernorm/gemm_add_relu_add_layernorm_welford.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <iostream>
...
...
@@ -78,7 +78,7 @@ int main(int argc, char* argv[])
[](
std
::
size_t
nRow
,
std
::
size_t
nCol
,
std
::
size_t
stride
,
auto
layout
)
{
using
Layout
=
decltype
(
layout
);
if
constexpr
(
std
::
is_same
<
Layout
,
ck
::
tensor_layout
::
gemm
::
RowMajor
>::
value
)
if
constexpr
(
std
::
is_same
<
Layout
,
Row
>::
value
)
{
return
(
nRow
-
1
)
*
stride
+
nCol
;
}
...
...
client_example/04_contraction/contraction_bilinear_fp32.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <numeric>
...
...
client_example/04_contraction/contraction_bilinear_fp64.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <numeric>
...
...
client_example/04_contraction/contraction_g1m2n3k1_add_xdl_fp16.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <numeric>
...
...
client_example/04_contraction/contraction_scale_fp32.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <numeric>
...
...
client_example/04_contraction/contraction_scale_fp64.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <numeric>
...
...
client_example/05_layernorm/layernorm2d_bwd_data.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
client_example/05_layernorm/layernorm2d_bwd_gamma_beta.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
client_example/05_layernorm/layernorm2d_fwd.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
client_example/05_layernorm/layernorm4d_fwd.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
...
...
client_example/06_softmax/softmax4d.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <functional>
#include <numeric>
...
...
client_example/07_grouped_convnd_fwd/grouped_conv1d_fwd.cpp
View file @
8ae8a558
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iomanip>
...
...
Prev
1
2
3
4
Next
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment