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
a245b8f3
Unverified
Commit
a245b8f3
authored
Nov 15, 2023
by
Illia Silin
Committed by
GitHub
Nov 15, 2023
Browse files
Merge branch 'develop' into lwpck-1026
parents
702228b0
f2398f61
Changes
247
Hide whitespace changes
Inline
Side-by-side
Showing
7 changed files
with
170 additions
and
2 deletions
+170
-2
test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp
test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp
+1
-1
test/normalization_fwd/test_layernorm2d_fwd_fp32.cpp
test/normalization_fwd/test_layernorm2d_fwd_fp32.cpp
+1
-1
test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp
test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp
+48
-0
test/transpose/CMakeLists.txt
test/transpose/CMakeLists.txt
+9
-0
test/transpose/test_transpose.cpp
test/transpose/test_transpose.cpp
+27
-0
test/transpose/test_transpose_ut_cases.inc
test/transpose/test_transpose_ut_cases.inc
+30
-0
test/transpose/test_transpose_util.hpp
test/transpose/test_transpose_util.hpp
+54
-0
No files found.
test/normalization/test_layernorm2d_fp16.cpp
→
test/normalization
_fwd
/test_layernorm2d_
fwd_
fp16.cpp
View file @
a245b8f3
...
...
@@ -2,7 +2,7 @@
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_layernorm_impl.hpp"
#include "profiler/profile_layernorm_
fwd_
impl.hpp"
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
...
...
test/normalization/test_layernorm2d_fp32.cpp
→
test/normalization
_fwd
/test_layernorm2d_
fwd_
fp32.cpp
View file @
a245b8f3
...
...
@@ -2,7 +2,7 @@
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_layernorm_impl.hpp"
#include "profiler/profile_layernorm_
fwd_
impl.hpp"
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
...
...
test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp
0 → 100644
View file @
a245b8f3
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_layernorm_fwd_impl.hpp"
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
ck
::
index_t
;
template
<
typename
Tuple
>
class
TestLayernorm4d
:
public
::
testing
::
Test
{
protected:
using
XDataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
GammaDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
BetaDataType
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
ComputeDataType
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
using
YDataType
=
std
::
tuple_element_t
<
4
,
Tuple
>
;
using
SaveMeanInvStdDataType
=
std
::
tuple_element_t
<
5
,
Tuple
>
;
void
Run
()
{
// [N, D], reduce D
std
::
vector
<
std
::
vector
<
ck
::
index_t
>>
lengths
=
{
{
1
,
1
,
1
,
1
},
{
7
,
7
,
7
,
7
},
{
256
,
16
,
16
,
8
}};
for
(
auto
length
:
lengths
)
{
bool
success
=
ck
::
profiler
::
profile_layernorm_impl
<
XDataType
,
GammaDataType
,
BetaDataType
,
ComputeDataType
,
YDataType
,
SaveMeanInvStdDataType
,
true
,
4
>
(
true
,
2
,
false
,
false
,
length
);
EXPECT_TRUE
(
success
);
}
}
};
using
KernelTypes
=
::
testing
::
Types
<
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType>
std
::
tuple
<
F16
,
F16
,
F16
,
F32
,
F16
,
F32
>>
;
TYPED_TEST_SUITE
(
TestLayernorm4d
,
KernelTypes
);
TYPED_TEST
(
TestLayernorm4d
,
Test_FP16
)
{
this
->
Run
();
}
test/transpose/CMakeLists.txt
0 → 100644
View file @
a245b8f3
list
(
APPEND gpu_list gfx908 gfx90a gfx940 gfx941 gfx942
)
set
(
target 0
)
foreach
(
gpu IN LISTS GPU_TARGETS
)
if
(
gpu IN_LIST gpu_list AND target EQUAL 0
)
add_gtest_executable
(
test_transpose test_transpose.cpp
)
target_link_libraries
(
test_transpose PRIVATE utility device_transpose_instance
)
set
(
target 1
)
endif
()
endforeach
()
test/transpose/test_transpose.cpp
0 → 100644
View file @
a245b8f3
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include <tuple>
#include "gtest/gtest.h"
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "test_transpose_util.hpp"
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
template
<
typename
Tuple
>
class
TestTranspose
:
public
::
testing
::
Test
{
};
// clang-format off
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
F16
>
,
std
::
tuple
<
F32
,
F32
>
>
;
// clang-format on
TYPED_TEST_SUITE
(
TestTranspose
,
KernelTypes
);
//#include "test_transpose_ut_cases.inc"
test/transpose/test_transpose_ut_cases.inc
0 → 100644
View file @
a245b8f3
#pragma once
TYPED_TEST
(
TestTranspose
,
Test1
)
{
// for 16, 8, 16, 32, 8
std
::
vector
<
int
>
Ms
{
1
,
2
,
3
,
4
,
5
,
6
};
std
::
vector
<
index_t
>
lengths
{
16
,
8
,
16
,
32
,
8
};
/**constexpr int N = 16;
constexpr int C = 8;
constexpr int D = 16;
constexpr int H = 32;
constexpr int W = 8;**/
this
->
Run
();
}
TYPED_TEST
(
TestTranpose
,
Test2
)
{
std
::
vector
<
int
>
Ms
{
127
,
255
,
312
,
799
,
1573
};
std
::
vector
<
index_t
>
lengths
{
16
,
8
,
16
,
32
,
16
};
/**constexpr int N = 16;
constexpr int C = 8;
constexpr int D = 16;
constexpr int H = 32;
constexpr int W = 8;**/
this
->
Run
();
}
test/transpose/test_transpose_util.hpp
0 → 100644
View file @
a245b8f3
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include <string>
#include <sstream>
#include <tuple>
#include <vector>
#include <gtest/gtest.h>
#include "ck/ck.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "include/ck/utility/data_type.hpp"
#include "profiler/profile_transpose_impl.hpp"
namespace
ck
{
namespace
test
{
template
<
typename
Tuple
>
class
TestTranspose
:
public
testing
::
Test
{
using
F32
=
float
;
protected:
using
ADataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
BDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
public:
static
constexpr
bool
verify_
=
true
;
static
constexpr
int
init_method_
=
1
;
// decimal value initialization
static
constexpr
bool
log_
=
false
;
static
constexpr
bool
bench_
=
false
;
// measure kernel performance
std
::
vector
<
std
::
vector
<
index_t
>>
lengths_
=
{{
16
,
32
,
16
,
32
,
16
},
{
16
,
8
,
16
,
32
,
8
}};
void
Run
()
{
for
(
auto
length
:
this
->
lengths_
)
{
this
->
RunSingle
(
length
);
}
}
void
RunSingle
()
{
bool
pass
=
ck
::
profiler
::
profile_transpose_impl
<
ADataType
,
BDataType
,
5
>
(
verify_
,
init_method_
,
log_
,
bench_
,
lengths_
);
EXPECT_TRUE
(
pass
);
}
};
}
// namespace test
}
// namespace ck
Prev
1
…
9
10
11
12
13
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