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
c14a0514
Commit
c14a0514
authored
Jan 31, 2024
by
Jakub Piasecki
Browse files
add tests for mixed precision gemms
parent
05fd7ff8
Changes
6
Show whitespace changes
Inline
Side-by-side
Showing
6 changed files
with
205 additions
and
0 deletions
+205
-0
test/CMakeLists.txt
test/CMakeLists.txt
+1
-0
test/gemm_add/CMakeLists.txt
test/gemm_add/CMakeLists.txt
+11
-0
test/gemm_add/test_gemm_add.hpp
test/gemm_add/test_gemm_add.hpp
+69
-0
test/gemm_add/test_gemm_add_fastgelu.cpp
test/gemm_add/test_gemm_add_fastgelu.cpp
+42
-0
test/gemm_add/test_gemm_add_relu.cpp
test/gemm_add/test_gemm_add_relu.cpp
+41
-0
test/gemm_add/test_gemm_add_silu.cpp
test/gemm_add/test_gemm_add_silu.cpp
+41
-0
No files found.
test/CMakeLists.txt
View file @
c14a0514
...
@@ -122,6 +122,7 @@ add_subdirectory(space_filling_curve)
...
@@ -122,6 +122,7 @@ add_subdirectory(space_filling_curve)
add_subdirectory
(
conv_util
)
add_subdirectory
(
conv_util
)
add_subdirectory
(
reference_conv_fwd
)
add_subdirectory
(
reference_conv_fwd
)
add_subdirectory
(
gemm
)
add_subdirectory
(
gemm
)
add_subdirectory
(
gemm_add
)
add_subdirectory
(
gemm_layernorm
)
add_subdirectory
(
gemm_layernorm
)
add_subdirectory
(
gemm_split_k
)
add_subdirectory
(
gemm_split_k
)
add_subdirectory
(
gemm_reduce
)
add_subdirectory
(
gemm_reduce
)
...
...
test/gemm_add/CMakeLists.txt
0 → 100644
View file @
c14a0514
add_gtest_executable
(
test_gemm_add test_gemm_add.hpp
)
target_link_libraries
(
test_gemm_add PRIVATE utility device_gemm_add_instance
)
add_gtest_executable
(
test_gemm_add_relu test_gemm_add_relu.cpp
)
target_link_libraries
(
test_gemm_add_relu PRIVATE utility device_gemm_add_instance device_gemm_add_relu_instance
)
add_gtest_executable
(
test_gemm_add_silu test_gemm_add_silu.cpp
)
target_link_libraries
(
test_gemm_add_silu PRIVATE utility device_gemm_add_instance device_gemm_add_silu_instance
)
add_gtest_executable
(
test_gemm_add_fastgelu test_gemm_add_fastgelu.cpp
)
target_link_libraries
(
test_gemm_add_fastgelu PRIVATE utility device_gemm_add_instance device_gemm_add_fastgelu_instance
)
test/gemm_add/test_gemm_add.hpp
0 → 100644
View file @
c14a0514
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "ck/ck.hpp"
#include "profiler/profile_gemm_add_impl.hpp"
using
Row
=
ck
::
tensor_layout
::
gemm
::
RowMajor
;
using
Col
=
ck
::
tensor_layout
::
gemm
::
ColumnMajor
;
using
I8
=
int8_t
;
using
BF16
=
ck
::
bhalf_t
;
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
template
<
typename
Tuple
>
class
TestGemmAdd
:
public
::
testing
::
Test
{
protected:
using
ADataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
BDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
AccDataType
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
D0DataType
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
using
EDataType
=
std
::
tuple_element_t
<
4
,
Tuple
>
;
using
ALayout
=
std
::
tuple_element_t
<
5
,
Tuple
>
;
using
BLayout
=
std
::
tuple_element_t
<
6
,
Tuple
>
;
using
D0Layout
=
std
::
tuple_element_t
<
7
,
Tuple
>
;
using
ELayout
=
std
::
tuple_element_t
<
8
,
Tuple
>
;
constexpr
static
auto
ProfileGemmAddImpl
=
ck
::
profiler
::
profile_gemm_add_impl
<
ADataType
,
BDataType
,
AccDataType
,
D0DataType
,
EDataType
,
ALayout
,
BLayout
,
D0Layout
,
ELayout
>
;
virtual
decltype
(
ProfileGemmAddImpl
)
GetImpl
()
{
return
ProfileGemmAddImpl
;
}
void
Run
()
{
std
::
vector
<
std
::
vector
<
ck
::
index_t
>>
lengths
=
{
{
16
,
32
,
64
},
{
2048
,
4096
,
8192
},
{
2048
,
4096
}};
for
(
auto
length
:
lengths
)
{
int
M
=
length
[
0
];
int
N
=
length
[
1
];
int
K
=
length
[
2
];
int
StrideA
=
ck
::
is_same_v
<
ALayout
,
Row
>
?
K
:
M
;
int
StrideB
=
ck
::
is_same_v
<
BLayout
,
Row
>
?
N
:
K
;
int
StrideD0
=
ck
::
is_same_v
<
D0Layout
,
Row
>
?
N
:
M
;
int
StrideE
=
ck
::
is_same_v
<
ELayout
,
Row
>
?
N
:
M
;
bool
success
=
GetImpl
()(
true
,
1
,
false
,
false
,
M
,
N
,
K
,
StrideA
,
StrideB
,
StrideD0
,
StrideE
);
EXPECT_TRUE
(
success
);
}
}
};
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
I8
,
F32
,
F16
,
F16
,
Row
,
Row
,
Row
,
Row
>
,
std
::
tuple
<
BF16
,
I8
,
F32
,
BF16
,
BF16
,
Row
,
Row
,
Row
,
Row
>>
;
TYPED_TEST_SUITE
(
TestGemmAdd
,
KernelTypes
);
TYPED_TEST
(
TestGemmAdd
,
Test_BF16FP16_INT8
)
{
this
->
Run
();
}
test/gemm_add/test_gemm_add_fastgelu.cpp
0 → 100644
View file @
c14a0514
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "ck/ck.hpp"
#include "profiler/profile_gemm_add_fastgelu_impl.hpp"
#include "test_gemm_add.hpp"
template
<
typename
Tuple
>
class
TestGemmAddFastgelu
:
public
TestGemmAdd
<
Tuple
>
{
private:
using
ADataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
BDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
AccDataType
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
D0DataType
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
using
EDataType
=
std
::
tuple_element_t
<
4
,
Tuple
>
;
using
ALayout
=
std
::
tuple_element_t
<
5
,
Tuple
>
;
using
BLayout
=
std
::
tuple_element_t
<
6
,
Tuple
>
;
using
D0Layout
=
std
::
tuple_element_t
<
7
,
Tuple
>
;
using
ELayout
=
std
::
tuple_element_t
<
8
,
Tuple
>
;
constexpr
static
auto
ProfileGemmAddFastgeluImpl
=
ck
::
profiler
::
profile_gemm_add_fastgelu_impl
<
ADataType
,
BDataType
,
AccDataType
,
D0DataType
,
EDataType
,
ALayout
,
BLayout
,
D0Layout
,
ELayout
>
;
decltype
(
ProfileGemmAddFastgeluImpl
)
GetImpl
()
override
{
return
ProfileGemmAddFastgeluImpl
;
}
};
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
F16
,
F32
,
F16
,
F16
,
Row
,
Row
,
Row
,
Row
>
,
std
::
tuple
<
F16
,
I8
,
F32
,
F16
,
F16
,
Row
,
Row
,
Row
,
Row
>
,
std
::
tuple
<
BF16
,
I8
,
F32
,
BF16
,
BF16
,
Row
,
Row
,
Row
,
Row
>>
;
TYPED_TEST_SUITE
(
TestGemmAddFastgelu
,
KernelTypes
);
TYPED_TEST
(
TestGemmAddFastgelu
,
Test_BF16FP16
)
{
this
->
Run
();
}
test/gemm_add/test_gemm_add_relu.cpp
0 → 100644
View file @
c14a0514
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "ck/ck.hpp"
#include "profiler/profile_gemm_add_relu_impl.hpp"
#include "test_gemm_add.hpp"
template
<
typename
Tuple
>
class
TestGemmAddRelu
:
public
TestGemmAdd
<
Tuple
>
{
private:
using
ADataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
BDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
AccDataType
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
D0DataType
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
using
EDataType
=
std
::
tuple_element_t
<
4
,
Tuple
>
;
using
ALayout
=
std
::
tuple_element_t
<
5
,
Tuple
>
;
using
BLayout
=
std
::
tuple_element_t
<
6
,
Tuple
>
;
using
D0Layout
=
std
::
tuple_element_t
<
7
,
Tuple
>
;
using
ELayout
=
std
::
tuple_element_t
<
8
,
Tuple
>
;
constexpr
static
auto
ProfileGemmAddReluImpl
=
ck
::
profiler
::
profile_gemm_add_relu_impl
<
ADataType
,
BDataType
,
AccDataType
,
D0DataType
,
EDataType
,
ALayout
,
BLayout
,
D0Layout
,
ELayout
>
;
decltype
(
ProfileGemmAddReluImpl
)
GetImpl
()
override
{
return
ProfileGemmAddReluImpl
;
}
};
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
I8
,
F32
,
F16
,
F16
,
Row
,
Row
,
Row
,
Row
>
,
std
::
tuple
<
BF16
,
I8
,
F32
,
BF16
,
BF16
,
Row
,
Row
,
Row
,
Row
>>
;
TYPED_TEST_SUITE
(
TestGemmAddRelu
,
KernelTypes
);
TYPED_TEST
(
TestGemmAddRelu
,
Test_BF16FP16_INT8
)
{
this
->
Run
();
}
test/gemm_add/test_gemm_add_silu.cpp
0 → 100644
View file @
c14a0514
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "ck/ck.hpp"
#include "profiler/profile_gemm_add_silu_impl.hpp"
#include "test_gemm_add.hpp"
template
<
typename
Tuple
>
class
TestGemmAddSilu
:
public
TestGemmAdd
<
Tuple
>
{
private:
using
ADataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
BDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
AccDataType
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
D0DataType
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
using
EDataType
=
std
::
tuple_element_t
<
4
,
Tuple
>
;
using
ALayout
=
std
::
tuple_element_t
<
5
,
Tuple
>
;
using
BLayout
=
std
::
tuple_element_t
<
6
,
Tuple
>
;
using
D0Layout
=
std
::
tuple_element_t
<
7
,
Tuple
>
;
using
ELayout
=
std
::
tuple_element_t
<
8
,
Tuple
>
;
constexpr
static
auto
ProfileGemmAddSiluImpl
=
ck
::
profiler
::
profile_gemm_add_silu_impl
<
ADataType
,
BDataType
,
AccDataType
,
D0DataType
,
EDataType
,
ALayout
,
BLayout
,
D0Layout
,
ELayout
>
;
decltype
(
ProfileGemmAddSiluImpl
)
GetImpl
()
override
{
return
ProfileGemmAddSiluImpl
;
}
};
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
I8
,
F32
,
F16
,
F16
,
Row
,
Row
,
Row
,
Row
>
,
std
::
tuple
<
BF16
,
I8
,
F32
,
BF16
,
BF16
,
Row
,
Row
,
Row
,
Row
>>
;
TYPED_TEST_SUITE
(
TestGemmAddSilu
,
KernelTypes
);
TYPED_TEST
(
TestGemmAddSilu
,
Test_BF16FP16_INT8
)
{
this
->
Run
();
}
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