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
Commits
40f4019d
Commit
40f4019d
authored
Jan 03, 2023
by
rocking
Browse files
Add test
parent
ddf71f54
Changes
4
Hide whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
91 additions
and
8 deletions
+91
-8
test/CMakeLists.txt
test/CMakeLists.txt
+2
-1
test/gemm_layernorm/CMakeLists.txt
test/gemm_layernorm/CMakeLists.txt
+7
-0
test/gemm_layernorm/test_gemm_add_relu_add_layernorm_fp16.cpp
.../gemm_layernorm/test_gemm_add_relu_add_layernorm_fp16.cpp
+76
-0
test/normalization/CMakeLists.txt
test/normalization/CMakeLists.txt
+6
-7
No files found.
test/CMakeLists.txt
View file @
40f4019d
...
@@ -27,7 +27,7 @@ function(add_gtest_executable TEST_NAME)
...
@@ -27,7 +27,7 @@ function(add_gtest_executable TEST_NAME)
# suppress gtest warnings
# suppress gtest warnings
target_compile_options
(
${
TEST_NAME
}
PRIVATE -Wno-global-constructors -Wno-undef
)
target_compile_options
(
${
TEST_NAME
}
PRIVATE -Wno-global-constructors -Wno-undef
)
target_link_libraries
(
${
TEST_NAME
}
PRIVATE gtest_main
)
target_link_libraries
(
${
TEST_NAME
}
PRIVATE gtest_main
)
add_test
(
NAME
${
TEST_NAME
}
COMMAND $<TARGET_FILE:
${
TEST_NAME
}
>
)
add_test
(
NAME
${
TEST_NAME
}
COMMAND $<TARGET_FILE:
${
TEST_NAME
}
>
)
rocm_install
(
TARGETS
${
TEST_NAME
}
COMPONENT tests
)
rocm_install
(
TARGETS
${
TEST_NAME
}
COMPONENT tests
)
endfunction
(
add_gtest_executable TEST_NAME
)
endfunction
(
add_gtest_executable TEST_NAME
)
...
@@ -36,6 +36,7 @@ add_subdirectory(space_filling_curve)
...
@@ -36,6 +36,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_layernorm
)
add_subdirectory
(
gemm_split_k
)
add_subdirectory
(
gemm_split_k
)
add_subdirectory
(
gemm_reduce
)
add_subdirectory
(
gemm_reduce
)
add_subdirectory
(
batched_gemm
)
add_subdirectory
(
batched_gemm
)
...
...
test/gemm_layernorm/CMakeLists.txt
0 → 100644
View file @
40f4019d
add_custom_target
(
test_gemm_layernorm
)
add_gtest_executable
(
test_gemm_add_relu_add_layernorm_fp16 test_gemm_add_relu_add_layernorm_fp16.cpp
)
target_link_libraries
(
test_gemm_add_relu_add_layernorm_fp16 PRIVATE utility device_gemm_add_relu_add_layernorm_instance
)
add_dependencies
(
test_gemm_layernorm test_gemm_add_relu_add_layernorm_fp16
)
test/gemm_layernorm/test_gemm_add_relu_add_layernorm_fp16.cpp
0 → 100644
View file @
40f4019d
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_gemm_add_relu_add_layernorm_impl.hpp"
using
Row
=
ck
::
tensor_layout
::
gemm
::
RowMajor
;
using
Col
=
ck
::
tensor_layout
::
gemm
::
ColumnMajor
;
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
ck
::
index_t
;
template
<
typename
Tuple
>
class
TestGemmAddReluAddLayernorm
:
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
D1DataType
=
std
::
tuple_element_t
<
4
,
Tuple
>
;
using
EMeanVarDataType
=
std
::
tuple_element_t
<
5
,
Tuple
>
;
using
GammaDataType
=
std
::
tuple_element_t
<
6
,
Tuple
>
;
using
BetaDataType
=
std
::
tuple_element_t
<
7
,
Tuple
>
;
using
HDataType
=
std
::
tuple_element_t
<
8
,
Tuple
>
;
using
ALayout
=
std
::
tuple_element_t
<
9
,
Tuple
>
;
using
BLayout
=
std
::
tuple_element_t
<
10
,
Tuple
>
;
using
D0Layout
=
std
::
tuple_element_t
<
11
,
Tuple
>
;
using
D1Layout
=
std
::
tuple_element_t
<
12
,
Tuple
>
;
using
HLayout
=
std
::
tuple_element_t
<
13
,
Tuple
>
;
void
Run
()
{
std
::
vector
<
std
::
vector
<
ck
::
index_t
>>
lengths
=
{{
1024
,
1024
,
1024
}};
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
=
0
;
int
StrideD1
=
ck
::
is_same_v
<
D1Layout
,
Row
>
?
N
:
M
;
int
StrideH
=
ck
::
is_same_v
<
HLayout
,
Row
>
?
N
:
M
;
bool
success
=
ck
::
profiler
::
profile_gemm_add_relu_add_layernorm_impl
<
ADataType
,
BDataType
,
AccDataType
,
D0DataType
,
D1DataType
,
EMeanVarDataType
,
GammaDataType
,
BetaDataType
,
HDataType
,
ALayout
,
BLayout
,
D0Layout
,
D1Layout
,
HLayout
>
(
true
,
1
,
false
,
false
,
M
,
N
,
K
,
StrideA
,
StrideB
,
StrideD0
,
StrideD1
,
StrideH
);
EXPECT_TRUE
(
success
);
}
}
};
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
F16
,
F32
,
F16
,
F16
,
F16
,
F16
,
F16
,
F16
,
Row
,
Row
,
Row
,
Row
,
Row
>
,
std
::
tuple
<
F16
,
F16
,
F32
,
F16
,
F16
,
F16
,
F16
,
F16
,
F16
,
Row
,
Col
,
Row
,
Row
,
Row
>
,
std
::
tuple
<
F16
,
F16
,
F32
,
F16
,
F16
,
F16
,
F16
,
F16
,
F16
,
Col
,
Row
,
Row
,
Row
,
Row
>
,
std
::
tuple
<
F16
,
F16
,
F32
,
F16
,
F16
,
F16
,
F16
,
F16
,
F16
,
Col
,
Col
,
Row
,
Row
,
Row
>>
;
TYPED_TEST_SUITE
(
TestGemmAddReluAddLayernorm
,
KernelTypes
);
TYPED_TEST
(
TestGemmAddReluAddLayernorm
,
Test_FP16
)
{
this
->
Run
();
}
test/normalization/CMakeLists.txt
View file @
40f4019d
add_custom_target
(
test_
layernorm
)
add_custom_target
(
test_
normalization
)
add_gtest_executable
(
test_layernorm2d_fp32 test_layernorm2d_fp32.cpp
)
add_gtest_executable
(
test_layernorm2d_fp32 test_layernorm2d_fp32.cpp
)
add_gtest_executable
(
test_layernorm2d_fp16 test_layernorm2d_fp16.cpp
)
add_gtest_executable
(
test_layernorm2d_fp16 test_layernorm2d_fp16.cpp
)
add_gtest_executable
(
test_groupnorm_fp16 test_groupnorm_fp16.cpp
)
add_gtest_executable
(
test_groupnorm_fp16 test_groupnorm_fp16.cpp
)
add_gtest_executable
(
test_groupnorm_fp32 test_groupnorm_fp32.cpp
)
add_gtest_executable
(
test_groupnorm_fp32 test_groupnorm_fp32.cpp
)
target_link_libraries
(
test_layernorm2d_fp32 PRIVATE utility device_normalization_instance
)
target_link_libraries
(
test_layernorm2d_fp32 PRIVATE utility device_normalization_instance
)
target_link_libraries
(
test_layernorm2d_fp16 PRIVATE utility device_normalization_instance
)
target_link_libraries
(
test_layernorm2d_fp16 PRIVATE utility device_normalization_instance
)
target_link_libraries
(
test_groupnorm_fp16 PRIVATE utility device_normalization_instance
)
target_link_libraries
(
test_groupnorm_fp16 PRIVATE utility device_normalization_instance
)
target_link_libraries
(
test_groupnorm_fp32 PRIVATE utility device_normalization_instance
)
target_link_libraries
(
test_groupnorm_fp32 PRIVATE utility device_normalization_instance
)
add_dependencies
(
test_
layernorm
test_layernorm2d_fp32
)
add_dependencies
(
test_
normalization
test_layernorm2d_fp32
)
add_dependencies
(
test_
layernorm
test_layernorm2d_fp16
)
add_dependencies
(
test_
normalization
test_layernorm2d_fp16
)
add_dependencies
(
test_
layernorm
test_groupnorm_fp16
)
add_dependencies
(
test_
normalization
test_groupnorm_fp16
)
add_dependencies
(
test_
layernorm
test_groupnorm_fp32
)
add_dependencies
(
test_
normalization
test_groupnorm_fp32
)
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