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
7df98b04
Commit
7df98b04
authored
Dec 01, 2023
by
rocking
Browse files
Change save mean var type from f32 to f16 in f16 mode
parent
de1cb1ed
Changes
15
Hide whitespace changes
Inline
Side-by-side
Showing
15 changed files
with
58 additions
and
58 deletions
+58
-58
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
library/include/ck/library/tensor_operation_instance/gpu/normalization_fwd.hpp
...brary/tensor_operation_instance/gpu/normalization_fwd.hpp
+4
-4
library/include/ck/library/tensor_operation_instance/gpu/normalization_fwd_swish.hpp
...tensor_operation_instance/gpu/normalization_fwd_swish.hpp
+2
-2
library/src/tensor_operation_instance/gpu/normalization_fwd/device_groupnorm_fwd_f16_instance.cpp
...u/normalization_fwd/device_groupnorm_fwd_f16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/normalization_fwd/device_groupnorm_fwd_swish_f16_instance.cpp
...alization_fwd/device_groupnorm_fwd_swish_f16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/normalization_fwd/device_layernorm2d_fwd_f16_instance.cpp
...normalization_fwd/device_layernorm2d_fwd_f16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/normalization_fwd/device_layernorm4d_fwd_f16_instance.cpp
...normalization_fwd/device_layernorm4d_fwd_f16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/normalization_fwd/normalization_fwd_instance_common.hpp
...u/normalization_fwd/normalization_fwd_instance_common.hpp
+37
-37
profiler/src/profile_groupnorm_fwd.cpp
profiler/src/profile_groupnorm_fwd.cpp
+1
-1
profiler/src/profile_layernorm_fwd.cpp
profiler/src/profile_layernorm_fwd.cpp
+1
-1
test/normalization_fwd/test_groupnorm_fwd_fp16.cpp
test/normalization_fwd/test_groupnorm_fwd_fp16.cpp
+2
-2
test/normalization_fwd/test_groupnorm_fwd_fp32.cpp
test/normalization_fwd/test_groupnorm_fwd_fp32.cpp
+1
-1
test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp
test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp
+2
-2
test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp
test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp
+2
-2
No files found.
client_example/05_layernorm/layernorm2d_fwd.cpp
View file @
7df98b04
...
@@ -16,7 +16,7 @@ using XDataType = ck::half_t;
...
@@ -16,7 +16,7 @@ using XDataType = ck::half_t;
using
GammaDataType
=
ck
::
half_t
;
using
GammaDataType
=
ck
::
half_t
;
using
BetaDataType
=
ck
::
half_t
;
using
BetaDataType
=
ck
::
half_t
;
using
YDataType
=
ck
::
half_t
;
using
YDataType
=
ck
::
half_t
;
using
SaveMeanInvStdDataType
=
floa
t
;
using
SaveMeanInvStdDataType
=
ck
::
half_
t
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
#define SAVE_MEAN_INV_STD
#define SAVE_MEAN_INV_STD
...
...
client_example/05_layernorm/layernorm4d_fwd.cpp
View file @
7df98b04
...
@@ -16,7 +16,7 @@ using XDataType = ck::half_t;
...
@@ -16,7 +16,7 @@ using XDataType = ck::half_t;
using
GammaDataType
=
ck
::
half_t
;
using
GammaDataType
=
ck
::
half_t
;
using
BetaDataType
=
ck
::
half_t
;
using
BetaDataType
=
ck
::
half_t
;
using
YDataType
=
ck
::
half_t
;
using
YDataType
=
ck
::
half_t
;
using
SaveMeanInvStdDataType
=
floa
t
;
using
SaveMeanInvStdDataType
=
ck
::
half_
t
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
#define SAVE_MEAN_INV_STD
#define SAVE_MEAN_INV_STD
...
...
library/include/ck/library/tensor_operation_instance/gpu/normalization_fwd.hpp
View file @
7df98b04
...
@@ -20,15 +20,15 @@ namespace instance {
...
@@ -20,15 +20,15 @@ namespace instance {
// FP16
// FP16
void
add_device_normalization_fwd_rank_2_1_f16_instances
(
void
add_device_normalization_fwd_rank_2_1_f16_instances
(
std
::
vector
<
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
32
,
PassThrough
,
2
,
1
>>>&
);
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
16
,
PassThrough
,
2
,
1
>>>&
);
void
add_device_normalization_fwd_rank_4_3_f16_instances
(
void
add_device_normalization_fwd_rank_4_3_f16_instances
(
std
::
vector
<
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
32
,
PassThrough
,
4
,
3
>>>&
);
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
16
,
PassThrough
,
4
,
3
>>>&
);
void
add_device_normalization_fwd_rank_5_3_f16_instances
(
void
add_device_normalization_fwd_rank_5_3_f16_instances
(
std
::
vector
<
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
32
,
PassThrough
,
5
,
3
>>>&
);
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
16
,
PassThrough
,
5
,
3
>>>&
);
#endif
#endif
#ifdef CK_ENABLE_FP32
#ifdef CK_ENABLE_FP32
// FP32
// FP32
...
@@ -76,7 +76,7 @@ struct DeviceOperationInstanceFactory<ck::tensor_operation::device::DeviceNormal
...
@@ -76,7 +76,7 @@ struct DeviceOperationInstanceFactory<ck::tensor_operation::device::DeviceNormal
#ifdef CK_ENABLE_FP16
#ifdef CK_ENABLE_FP16
if
constexpr
(
is_same_v
<
XDataType
,
F16
>
&&
is_same_v
<
GammaDataType
,
F16
>
&&
if
constexpr
(
is_same_v
<
XDataType
,
F16
>
&&
is_same_v
<
GammaDataType
,
F16
>
&&
is_same_v
<
BetaDataType
,
F16
>
&&
is_same_v
<
YDataType
,
F16
>
&&
is_same_v
<
BetaDataType
,
F16
>
&&
is_same_v
<
YDataType
,
F16
>
&&
is_same_v
<
SaveMeanInvStdDataType
,
F
32
>
)
is_same_v
<
SaveMeanInvStdDataType
,
F
16
>
)
{
{
if
constexpr
(
Rank
==
2
&&
NumReduceDim
==
1
)
if
constexpr
(
Rank
==
2
&&
NumReduceDim
==
1
)
{
{
...
...
library/include/ck/library/tensor_operation_instance/gpu/normalization_fwd_swish.hpp
View file @
7df98b04
...
@@ -19,7 +19,7 @@ namespace instance {
...
@@ -19,7 +19,7 @@ namespace instance {
// FP16
// FP16
void
add_device_normalization_fwd_rank_5_3_swish_f16_instances
(
void
add_device_normalization_fwd_rank_5_3_swish_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
32
,
Swish
,
5
,
3
>>>&
);
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
16
,
Swish
,
5
,
3
>>>&
);
// FP32
// FP32
void
add_device_normalization_fwd_rank_5_3_swish_f32_instances
(
void
add_device_normalization_fwd_rank_5_3_swish_f32_instances
(
...
@@ -61,7 +61,7 @@ struct DeviceOperationInstanceFactory<
...
@@ -61,7 +61,7 @@ struct DeviceOperationInstanceFactory<
if
constexpr
(
is_same_v
<
XDataType
,
F16
>
&&
is_same_v
<
GammaDataType
,
F16
>
&&
if
constexpr
(
is_same_v
<
XDataType
,
F16
>
&&
is_same_v
<
GammaDataType
,
F16
>
&&
is_same_v
<
BetaDataType
,
F16
>
&&
is_same_v
<
YDataType
,
F16
>
&&
is_same_v
<
BetaDataType
,
F16
>
&&
is_same_v
<
YDataType
,
F16
>
&&
is_same_v
<
SaveMeanInvStdDataType
,
F
32
>
)
is_same_v
<
SaveMeanInvStdDataType
,
F
16
>
)
{
{
if
constexpr
(
Rank
==
5
&&
NumReduceDim
==
3
)
if
constexpr
(
Rank
==
5
&&
NumReduceDim
==
3
)
{
{
...
...
library/src/tensor_operation_instance/gpu/normalization_fwd/device_groupnorm_fwd_f16_instance.cpp
View file @
7df98b04
...
@@ -11,7 +11,7 @@ namespace instance {
...
@@ -11,7 +11,7 @@ namespace instance {
using
Pass
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
Pass
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
void
add_device_normalization_fwd_rank_5_3_f16_instances
(
void
add_device_normalization_fwd_rank_5_3_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
32
,
Pass
,
5
,
3
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
16
,
Pass
,
5
,
3
>>>&
instances
)
instances
)
{
{
add_device_operation_instances
(
instances
,
add_device_operation_instances
(
instances
,
...
...
library/src/tensor_operation_instance/gpu/normalization_fwd/device_groupnorm_fwd_swish_f16_instance.cpp
View file @
7df98b04
...
@@ -11,7 +11,7 @@ namespace instance {
...
@@ -11,7 +11,7 @@ namespace instance {
using
Swish
=
ck
::
tensor_operation
::
element_wise
::
Swish
;
using
Swish
=
ck
::
tensor_operation
::
element_wise
::
Swish
;
void
add_device_normalization_fwd_rank_5_3_swish_f16_instances
(
void
add_device_normalization_fwd_rank_5_3_swish_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
32
,
Swish
,
5
,
3
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
16
,
Swish
,
5
,
3
>>>&
instances
)
instances
)
{
{
add_device_operation_instances
(
instances
,
add_device_operation_instances
(
instances
,
...
...
library/src/tensor_operation_instance/gpu/normalization_fwd/device_layernorm2d_fwd_f16_instance.cpp
View file @
7df98b04
...
@@ -11,7 +11,7 @@ namespace instance {
...
@@ -11,7 +11,7 @@ namespace instance {
using
Pass
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
Pass
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
void
add_device_normalization_fwd_rank_2_1_f16_instances
(
void
add_device_normalization_fwd_rank_2_1_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
32
,
Pass
,
2
,
1
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
16
,
Pass
,
2
,
1
>>>&
instances
)
instances
)
{
{
add_device_operation_instances
(
instances
,
add_device_operation_instances
(
instances
,
...
...
library/src/tensor_operation_instance/gpu/normalization_fwd/device_layernorm4d_fwd_f16_instance.cpp
View file @
7df98b04
...
@@ -11,7 +11,7 @@ namespace instance {
...
@@ -11,7 +11,7 @@ namespace instance {
using
Pass
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
Pass
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
void
add_device_normalization_fwd_rank_4_3_f16_instances
(
void
add_device_normalization_fwd_rank_4_3_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
32
,
Pass
,
4
,
3
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceNormalizationFwd
<
F16
,
F16
,
F16
,
F16
,
F
16
,
Pass
,
4
,
3
>>>&
instances
)
instances
)
{
{
add_device_operation_instances
(
instances
,
add_device_operation_instances
(
instances
,
...
...
library/src/tensor_operation_instance/gpu/normalization_fwd/normalization_fwd_instance_common.hpp
View file @
7df98b04
...
@@ -23,24 +23,24 @@ using device_normalization_f16_instances =
...
@@ -23,24 +23,24 @@ using device_normalization_f16_instances =
// clang-format off
// clang-format off
std
::
tuple
<
std
::
tuple
<
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType, SaveMeanInvStdDataType, Rank, NumReduceDim, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, XYSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, BetaSrcVectorDim, BetaSrcVectorSize, YDstVectorSize, SaveMeanInvStdScalarPerVector>
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType, SaveMeanInvStdDataType, Rank, NumReduceDim, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, XYSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, BetaSrcVectorDim, BetaSrcVectorSize, YDstVectorSize, SaveMeanInvStdScalarPerVector>
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
2
,
1
,
2
,
1
,
2
,
1
,
2
,
2
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
2
,
1
,
2
,
1
,
2
,
1
,
2
,
2
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
4
,
1
,
4
,
1
,
4
,
1
,
4
,
4
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
4
,
1
,
4
,
1
,
4
,
1
,
4
,
4
,
1
>
,
// irregular size
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
64
,
1
,
64
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
64
,
1
,
64
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
32
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
32
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
2
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
2
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
2
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
2
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
32
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
32
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
// clang-format on
// clang-format on
>
;
>
;
...
@@ -49,31 +49,31 @@ using device_normalization_splitk_f16_instances =
...
@@ -49,31 +49,31 @@ using device_normalization_splitk_f16_instances =
// clang-format off
// clang-format off
std
::
tuple
<
std
::
tuple
<
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType, SaveMeanInvStdDataType, Rank, NumReduceDim, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, XYSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, BetaSrcVectorDim, BetaSrcVectorSize, YDstVectorSize, SaveMeanInvStdScalarPerVector>
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType, SaveMeanInvStdDataType, Rank, NumReduceDim, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, XYSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, BetaSrcVectorDim, BetaSrcVectorSize, YDstVectorSize, SaveMeanInvStdScalarPerVector>
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
2
,
1
,
2
,
1
,
2
,
1
,
2
,
2
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
2
,
1
,
2
,
1
,
2
,
1
,
2
,
2
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
4
,
1
,
4
,
1
,
4
,
1
,
4
,
4
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
4
,
1
,
4
,
1
,
4
,
1
,
4
,
4
,
1
>
,
// irregular size
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
64
,
1
,
64
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
64
,
1
,
64
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
32
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
128
,
1
,
128
,
1
,
32
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
2
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
2
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
2
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
2
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
32
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
256
,
1
,
256
,
1
,
32
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
512
,
1
,
512
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
8
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
,
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
DeviceNormalizationFwdSplitKImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
1024
,
1
,
1024
,
1
,
16
,
1
,
8
,
1
,
8
,
1
,
8
,
8
,
1
>
// clang-format on
// clang-format on
>
;
>
;
template
<
typename
OutElementwise
,
index_t
Rank
,
index_t
Reduce
>
template
<
typename
OutElementwise
,
index_t
Rank
,
index_t
Reduce
>
using
device_normalization_f16_generic_instance
=
std
::
tuple
<
using
device_normalization_f16_generic_instance
=
std
::
tuple
<
// clang-format off
// clang-format off
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
OutElementwise
,
Rank
,
Reduce
,
64
,
1
,
64
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
DeviceNormalizationFwdImpl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
OutElementwise
,
Rank
,
Reduce
,
64
,
1
,
64
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
,
1
>
// clang-format on
// clang-format on
>
;
>
;
...
...
profiler/src/profile_groupnorm_fwd.cpp
View file @
7df98b04
...
@@ -98,7 +98,7 @@ int profile_groupnorm(int argc, char* argv[])
...
@@ -98,7 +98,7 @@ int profile_groupnorm(int argc, char* argv[])
}
}
else
if
(
data_type
==
ck
::
DataTypeEnum
::
Half
)
else
if
(
data_type
==
ck
::
DataTypeEnum
::
Half
)
{
{
ck
::
profiler
::
profile_groupnorm_impl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
false
>
(
ck
::
profiler
::
profile_groupnorm_impl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
false
>
(
do_verification
,
init_method
,
do_log
,
time_kernel
,
length
);
do_verification
,
init_method
,
do_log
,
time_kernel
,
length
);
}
}
else
else
...
...
profiler/src/profile_layernorm_fwd.cpp
View file @
7df98b04
...
@@ -104,7 +104,7 @@ int profile_layernorm(int argc, char* argv[])
...
@@ -104,7 +104,7 @@ int profile_layernorm(int argc, char* argv[])
if
(
data_type
==
ck
::
DataTypeEnum
::
Half
)
if
(
data_type
==
ck
::
DataTypeEnum
::
Half
)
{
{
ck
::
profiler
::
profile_layernorm_impl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
,
false
,
rank
>
(
ck
::
profiler
::
profile_layernorm_impl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
,
false
,
rank
>
(
do_verification
,
init_method
,
do_log
,
time_kernel
,
length
);
do_verification
,
init_method
,
do_log
,
time_kernel
,
length
);
}
}
else
if
(
data_type
==
ck
::
DataTypeEnum
::
Float
)
else
if
(
data_type
==
ck
::
DataTypeEnum
::
Float
)
...
...
test/normalization_fwd/test_groupnorm_fwd_fp16.cpp
View file @
7df98b04
...
@@ -47,8 +47,8 @@ class TestGroupnorm : public ::testing::Test
...
@@ -47,8 +47,8 @@ class TestGroupnorm : public ::testing::Test
};
};
using
KernelTypes
=
::
testing
::
Types
<
using
KernelTypes
=
::
testing
::
Types
<
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType>
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType
, SaveMeanInvStdDataType
>
std
::
tuple
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
>>
;
std
::
tuple
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
>>
;
TYPED_TEST_SUITE
(
TestGroupnorm
,
KernelTypes
);
TYPED_TEST_SUITE
(
TestGroupnorm
,
KernelTypes
);
TYPED_TEST
(
TestGroupnorm
,
Test_FP16
)
{
this
->
Run
();
}
TYPED_TEST
(
TestGroupnorm
,
Test_FP16
)
{
this
->
Run
();
}
test/normalization_fwd/test_groupnorm_fwd_fp32.cpp
View file @
7df98b04
...
@@ -45,7 +45,7 @@ class TestGroupnorm : public ::testing::Test
...
@@ -45,7 +45,7 @@ class TestGroupnorm : public ::testing::Test
};
};
using
KernelTypes
=
::
testing
::
Types
<
using
KernelTypes
=
::
testing
::
Types
<
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType>
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType
, SaveMeanInvStdDataType
>
std
::
tuple
<
F32
,
F32
,
F32
,
F32
,
F32
,
F32
>>
;
std
::
tuple
<
F32
,
F32
,
F32
,
F32
,
F32
,
F32
>>
;
TYPED_TEST_SUITE
(
TestGroupnorm
,
KernelTypes
);
TYPED_TEST_SUITE
(
TestGroupnorm
,
KernelTypes
);
...
...
test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp
View file @
7df98b04
...
@@ -41,8 +41,8 @@ class TestLayernorm2d : public ::testing::Test
...
@@ -41,8 +41,8 @@ class TestLayernorm2d : public ::testing::Test
};
};
using
KernelTypes
=
::
testing
::
Types
<
using
KernelTypes
=
::
testing
::
Types
<
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType>
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType
, SaveMeanInvStdDataType
>
std
::
tuple
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
>>
;
std
::
tuple
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
>>
;
TYPED_TEST_SUITE
(
TestLayernorm2d
,
KernelTypes
);
TYPED_TEST_SUITE
(
TestLayernorm2d
,
KernelTypes
);
TYPED_TEST
(
TestLayernorm2d
,
Test_FP16
)
{
this
->
Run
();
}
TYPED_TEST
(
TestLayernorm2d
,
Test_FP16
)
{
this
->
Run
();
}
test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp
View file @
7df98b04
...
@@ -41,8 +41,8 @@ class TestLayernorm4d : public ::testing::Test
...
@@ -41,8 +41,8 @@ class TestLayernorm4d : public ::testing::Test
};
};
using
KernelTypes
=
::
testing
::
Types
<
using
KernelTypes
=
::
testing
::
Types
<
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType>
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType
, SaveMeanInvStdDataType
>
std
::
tuple
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
32
>>
;
std
::
tuple
<
F16
,
F16
,
F16
,
F32
,
F16
,
F
16
>>
;
TYPED_TEST_SUITE
(
TestLayernorm4d
,
KernelTypes
);
TYPED_TEST_SUITE
(
TestLayernorm4d
,
KernelTypes
);
TYPED_TEST
(
TestLayernorm4d
,
Test_FP16
)
{
this
->
Run
();
}
TYPED_TEST
(
TestLayernorm4d
,
Test_FP16
)
{
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