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
e9070031
Commit
e9070031
authored
Sep 16, 2022
by
rocking
Browse files
[What] Rename original layernorm into layernorm2d
[Why] Prepare to add groupnorm using layernorm5d
parent
58188d46
Changes
4
Hide whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
32 additions
and
34 deletions
+32
-34
test/layernorm/CMakeLists.txt
test/layernorm/CMakeLists.txt
+6
-6
test/layernorm/test_layernorm2d_fp16.cpp
test/layernorm/test_layernorm2d_fp16.cpp
+4
-4
test/layernorm/test_layernorm2d_fp32.cpp
test/layernorm/test_layernorm2d_fp32.cpp
+4
-4
test/layernorm/test_layernorm2d_util.hpp
test/layernorm/test_layernorm2d_util.hpp
+18
-20
No files found.
test/layernorm/CMakeLists.txt
View file @
e9070031
add_custom_target
(
test_layernorm
)
add_custom_target
(
test_layernorm
)
add_gtest_executable
(
test_layernorm_fp32 test_layernorm_fp32.cpp
)
add_gtest_executable
(
test_layernorm
2d
_fp32 test_layernorm
2d
_fp32.cpp
)
add_gtest_executable
(
test_layernorm_fp16 test_layernorm_fp16.cpp
)
add_gtest_executable
(
test_layernorm
2d
_fp16 test_layernorm
2d
_fp16.cpp
)
target_link_libraries
(
test_layernorm_fp32 PRIVATE utility
)
target_link_libraries
(
test_layernorm
2d
_fp32 PRIVATE utility
)
target_link_libraries
(
test_layernorm_fp16 PRIVATE utility
)
target_link_libraries
(
test_layernorm
2d
_fp16 PRIVATE utility
)
add_dependencies
(
test_layernorm test_layernorm_fp32
)
add_dependencies
(
test_layernorm test_layernorm
2d
_fp32
)
add_dependencies
(
test_layernorm test_layernorm_fp16
)
add_dependencies
(
test_layernorm test_layernorm
2d
_fp16
)
test/layernorm/test_layernorm_fp16.cpp
→
test/layernorm/test_layernorm
2d
_fp16.cpp
View file @
e9070031
...
@@ -2,13 +2,13 @@
...
@@ -2,13 +2,13 @@
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "gtest/gtest.h"
#include "test_layernorm_util.hpp"
#include "test_layernorm
2d
_util.hpp"
template
<
ck
::
index_t
N
>
template
<
ck
::
index_t
N
>
using
I
=
ck
::
Number
<
N
>
;
using
I
=
ck
::
Number
<
N
>
;
template
<
typename
Tuple
>
template
<
typename
Tuple
>
class
TestLayernormFP16
:
public
ck
::
TestLayernorm
<
Tuple
>
class
TestLayernorm
2d
FP16
:
public
ck
::
TestLayernorm
2d
<
Tuple
>
{
{
};
};
...
@@ -25,5 +25,5 @@ using KernelTypes = ::testing::Types<
...
@@ -25,5 +25,5 @@ using KernelTypes = ::testing::Types<
std
::
tuple
<
ck
::
half_t
,
ck
::
half_t
,
ck
::
half_t
,
float
,
ck
::
half_t
,
I
<
2
>
,
I
<
1
>
,
I
<
256
>
,
I
<
1
>
,
I
<
256
>
,
I
<
2
>
,
I
<
8
>
,
I
<
1
>
,
I
<
8
>
,
I
<
1
>
,
I
<
8
>
,
I
<
1
>
,
I
<
8
>
,
I
<
8
>>
std
::
tuple
<
ck
::
half_t
,
ck
::
half_t
,
ck
::
half_t
,
float
,
ck
::
half_t
,
I
<
2
>
,
I
<
1
>
,
I
<
256
>
,
I
<
1
>
,
I
<
256
>
,
I
<
2
>
,
I
<
8
>
,
I
<
1
>
,
I
<
8
>
,
I
<
1
>
,
I
<
8
>
,
I
<
1
>
,
I
<
8
>
,
I
<
8
>>
>
;
>
;
// clang-format on
// clang-format on
TYPED_TEST_SUITE
(
TestLayernormFP16
,
KernelTypes
);
TYPED_TEST_SUITE
(
TestLayernorm
2d
FP16
,
KernelTypes
);
TYPED_TEST
(
TestLayernormFP16
,
Test_FP16
)
{
this
->
Run
();
}
TYPED_TEST
(
TestLayernorm
2d
FP16
,
Test_FP16
)
{
this
->
Run
();
}
test/layernorm/test_layernorm_fp32.cpp
→
test/layernorm/test_layernorm
2d
_fp32.cpp
View file @
e9070031
...
@@ -2,13 +2,13 @@
...
@@ -2,13 +2,13 @@
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "gtest/gtest.h"
#include "test_layernorm_util.hpp"
#include "test_layernorm
2d
_util.hpp"
template
<
ck
::
index_t
N
>
template
<
ck
::
index_t
N
>
using
I
=
ck
::
Number
<
N
>
;
using
I
=
ck
::
Number
<
N
>
;
template
<
typename
Tuple
>
template
<
typename
Tuple
>
class
TestLayernormFP32
:
public
ck
::
TestLayernorm
<
Tuple
>
class
TestLayernorm
2d
FP32
:
public
ck
::
TestLayernorm
2d
<
Tuple
>
{
{
};
};
...
@@ -25,5 +25,5 @@ using KernelTypes = ::testing::Types<
...
@@ -25,5 +25,5 @@ using KernelTypes = ::testing::Types<
std
::
tuple
<
float
,
float
,
float
,
float
,
float
,
I
<
2
>
,
I
<
1
>
,
I
<
256
>
,
I
<
1
>
,
I
<
256
>
,
I
<
2
>
,
I
<
8
>
,
I
<
1
>
,
I
<
4
>
,
I
<
1
>
,
I
<
4
>
,
I
<
1
>
,
I
<
4
>
,
I
<
4
>>
std
::
tuple
<
float
,
float
,
float
,
float
,
float
,
I
<
2
>
,
I
<
1
>
,
I
<
256
>
,
I
<
1
>
,
I
<
256
>
,
I
<
2
>
,
I
<
8
>
,
I
<
1
>
,
I
<
4
>
,
I
<
1
>
,
I
<
4
>
,
I
<
1
>
,
I
<
4
>
,
I
<
4
>>
>
;
>
;
// clang-format on
// clang-format on
TYPED_TEST_SUITE
(
TestLayernormFP32
,
KernelTypes
);
TYPED_TEST_SUITE
(
TestLayernorm
2d
FP32
,
KernelTypes
);
TYPED_TEST
(
TestLayernormFP32
,
Test_FP32
)
{
this
->
Run
();
}
TYPED_TEST
(
TestLayernorm
2d
FP32
,
Test_FP32
)
{
this
->
Run
();
}
test/layernorm/test_layernorm_util.hpp
→
test/layernorm/test_layernorm
2d
_util.hpp
View file @
e9070031
...
@@ -31,7 +31,7 @@ std::string serialize_range(const Range& range)
...
@@ -31,7 +31,7 @@ std::string serialize_range(const Range& range)
}
}
template
<
typename
Tuple
>
template
<
typename
Tuple
>
class
TestLayernorm
:
public
::
testing
::
Test
class
TestLayernorm
2d
:
public
::
testing
::
Test
{
{
protected:
protected:
using
XDataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
XDataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
...
@@ -86,19 +86,18 @@ class TestLayernorm : public ::testing::Test
...
@@ -86,19 +86,18 @@ class TestLayernorm : public ::testing::Test
BetaSrcVectorSize
,
BetaSrcVectorSize
,
YDstVectorSize
>
;
YDstVectorSize
>
;
TestLayernorm
()
:
ref_instance_invoker_
(
ReferenceInstance
{}.
MakeInvoker
())
{}
TestLayernorm
2d
()
:
ref_instance_invoker_
(
ReferenceInstance
{}.
MakeInvoker
())
{}
void
RunSingle
(
std
::
vector
<
index_t
>
lengths
,
std
::
vector
<
index_t
>
reduceDims
)
void
RunSingle
(
const
std
::
vector
<
index_t
>&
lengths
,
const
std
::
vector
<
index_t
>&
reduceDims
,
const
std
::
vector
<
index_t
>&
GammaLength
,
const
std
::
vector
<
index_t
>&
GammaStride
,
const
std
::
vector
<
index_t
>&
BetaLength
,
const
std
::
vector
<
index_t
>&
BetaStride
)
{
{
std
::
vector
<
index_t
>
reduceLength
(
reduceDims
.
size
());
for
(
int
i
=
0
;
i
<
NumReduceDim
;
++
i
)
{
reduceLength
[
i
]
=
lengths
[
reduceDims
[
i
]];
}
Tensor
<
XDataType
>
x
(
lengths
);
Tensor
<
XDataType
>
x
(
lengths
);
Tensor
<
GammaDataType
>
gamma
(
reduce
Length
);
Tensor
<
GammaDataType
>
gamma
(
Gamma
Length
);
Tensor
<
BetaDataType
>
beta
(
reduce
Length
);
Tensor
<
BetaDataType
>
beta
(
Beta
Length
);
Tensor
<
YDataType
>
y
(
lengths
);
Tensor
<
YDataType
>
y
(
lengths
);
Tensor
<
YDataType
>
y_ref
(
lengths
);
Tensor
<
YDataType
>
y_ref
(
lengths
);
...
@@ -119,8 +118,8 @@ class TestLayernorm : public ::testing::Test
...
@@ -119,8 +118,8 @@ class TestLayernorm : public ::testing::Test
auto
argument_ptr
=
device_instance
.
MakeArgumentPointer
(
auto
argument_ptr
=
device_instance
.
MakeArgumentPointer
(
lengths
,
lengths
,
std
::
vector
<
ck
::
index_t
>
{
x
.
mDesc
.
GetStrides
().
begin
(),
x
.
mDesc
.
GetStrides
().
end
()},
std
::
vector
<
ck
::
index_t
>
{
x
.
mDesc
.
GetStrides
().
begin
(),
x
.
mDesc
.
GetStrides
().
end
()},
{
0
,
1
}
,
GammaStride
,
{
0
,
1
}
,
BetaStride
,
std
::
vector
<
ck
::
index_t
>
{
y
.
mDesc
.
GetStrides
().
begin
(),
y
.
mDesc
.
GetStrides
().
end
()},
std
::
vector
<
ck
::
index_t
>
{
y
.
mDesc
.
GetStrides
().
begin
(),
y
.
mDesc
.
GetStrides
().
end
()},
reduceDims
,
reduceDims
,
1e-4
,
1e-4
,
...
@@ -165,17 +164,16 @@ class TestLayernorm : public ::testing::Test
...
@@ -165,17 +164,16 @@ class TestLayernorm : public ::testing::Test
void
Run
()
void
Run
()
{
{
for
(
auto
length
:
this
->
lengths_
)
std
::
vector
<
std
::
vector
<
index_t
>>
lengths
=
{
{
4
,
256
},
{
8
,
511
},
{
9
,
1032
},
{
4
,
2048
},
{
1
,
8192
},
{
4000
,
2000
}};
for
(
auto
length
:
lengths
)
{
{
this
->
RunSingle
(
length
,
reduceDims_
[
0
]
);
this
->
RunSingle
(
length
,
{
1
},
{
length
[
1
]},
{
0
,
1
},
{
length
[
1
]},
{
0
,
1
}
);
}
}
}
}
std
::
vector
<
std
::
vector
<
index_t
>>
lengths_
=
{
{
4
,
256
},
{
8
,
511
},
{
9
,
1032
},
{
4
,
2048
},
{
1
,
8192
},
{
4000
,
2000
}};
std
::
vector
<
std
::
vector
<
index_t
>>
reduceDims_
=
{{
1
}};
typename
ReferenceInstance
::
Invoker
ref_instance_invoker_
;
typename
ReferenceInstance
::
Invoker
ref_instance_invoker_
;
};
};
}
// namespace ck
}
// namespace ck
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