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
5903efe7
Unverified
Commit
5903efe7
authored
Nov 16, 2023
by
arai713
Committed by
GitHub
Nov 16, 2023
Browse files
Merge branch 'develop' into transpose_5d
parents
2100ea4b
e1fa0091
Changes
231
Hide whitespace changes
Inline
Side-by-side
Showing
11 changed files
with
382 additions
and
26 deletions
+382
-26
test/CMakeLists.txt
test/CMakeLists.txt
+1
-1
test/grouped_convnd_fwd/CMakeLists.txt
test/grouped_convnd_fwd/CMakeLists.txt
+5
-0
test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_ab_interface.cpp
...convnd_fwd/test_grouped_convnd_fwd_multi_ab_interface.cpp
+235
-0
test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_d_interface_compatibility.cpp
...st_grouped_convnd_fwd_multi_d_interface_compatibility.cpp
+59
-0
test/normalization/CMakeLists.txt
test/normalization/CMakeLists.txt
+0
-21
test/normalization_fwd/CMakeLists.txt
test/normalization_fwd/CMakeLists.txt
+30
-0
test/normalization_fwd/test_groupnorm_fwd_fp16.cpp
test/normalization_fwd/test_groupnorm_fwd_fp16.cpp
+1
-1
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
+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
No files found.
test/CMakeLists.txt
View file @
5903efe7
...
...
@@ -139,7 +139,7 @@ add_subdirectory(grouped_convnd_fwd)
add_subdirectory
(
grouped_convnd_bwd_weight
)
add_subdirectory
(
block_to_ctile_map
)
add_subdirectory
(
softmax
)
add_subdirectory
(
normalization
)
add_subdirectory
(
normalization
_fwd
)
add_subdirectory
(
data_type
)
add_subdirectory
(
elementwise_normalization
)
add_subdirectory
(
batchnorm
)
...
...
test/grouped_convnd_fwd/CMakeLists.txt
View file @
5903efe7
add_gtest_executable
(
test_grouped_convnd_fwd test_grouped_convnd_fwd.cpp
)
target_link_libraries
(
test_grouped_convnd_fwd PRIVATE utility device_grouped_conv1d_fwd_instance device_grouped_conv2d_fwd_instance device_grouped_conv3d_fwd_instance
)
add_gtest_executable
(
test_grouped_convnd_fwd_multi_ab_interface test_grouped_convnd_fwd_multi_ab_interface.cpp
)
target_link_libraries
(
test_grouped_convnd_fwd_multi_ab_interface PRIVATE utility
)
add_gtest_executable
(
test_grouped_convnd_fwd_multi_d_interface_compatibility test_grouped_convnd_fwd_multi_d_interface_compatibility.cpp
)
target_link_libraries
(
test_grouped_convnd_fwd_multi_d_interface_compatibility PRIVATE utility device_grouped_conv3d_fwd_instance
)
test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_ab_interface.cpp
0 → 100644
View file @
5903efe7
// SPDX-License-Identifier: MIT
// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iostream>
#include <initializer_list>
#include <tuple>
#include <vector>
#include "ck/ck.hpp"
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp"
#include "ck/host_utility/device_prop.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
#include "ck/library/utility/algorithm.hpp"
#include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp"
#include <gtest/gtest.h>
template
<
ck
::
index_t
...
Is
>
using
S
=
ck
::
Sequence
<
Is
...
>
;
using
ScaleAdd
=
ck
::
tensor_operation
::
element_wise
::
ScaleAdd
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
template
<
typename
DataType
,
typename
InDataTypes
,
typename
WeiDataTypes
,
typename
InElementOp
,
typename
WeiElementOp
>
class
TestGroupedConvndFwdMultiABInterfaceBase
:
public
::
testing
::
Test
{
protected:
static
constexpr
ck
::
index_t
NDimSpatial
=
3
;
static
constexpr
ck
::
index_t
NumAs
=
2
;
static
constexpr
ck
::
index_t
NumBs
=
2
;
static
constexpr
auto
ConvSpec
=
ck
::
tensor_operation
::
device
::
ConvolutionForwardSpecialization
::
Default
;
static
constexpr
auto
GemmSpec
=
ck
::
tensor_operation
::
device
::
GemmSpecialization
::
MNKPadding
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
GNDHWC
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKZYXC
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
GNDHWK
;
using
OutElementOp
=
PassThrough
;
using
DeviceGroupedConvNDMultiABFwdInstance
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
<
NDimSpatial
,
InLayout
,
WeiLayout
,
ck
::
Tuple
<>
,
OutLayout
,
InDataTypes
,
WeiDataTypes
,
DataType
,
DataType
,
ck
::
Tuple
<>
,
DataType
,
InElementOp
,
WeiElementOp
,
OutElementOp
,
ConvSpec
,
// ConvForwardSpecialization
GemmSpec
,
// GemmSpecialization
1
,
//
256
,
// BlockSize
128
,
// MPerBlock
256
,
// NPerBlock
32
,
// KPerBlock
8
,
// AK1
8
,
// BK1
32
,
// MPerXdl
32
,
// NPerXdl
2
,
// MXdlPerWave
4
,
// NXdlPerWave
S
<
4
,
64
,
1
>
,
// ABlockTransferThreadClusterLengths_AK0_M_AK1
S
<
1
,
0
,
2
>
,
// ABlockTransferThreadClusterArrangeOrder
S
<
1
,
0
,
2
>
,
// ABlockTransferSrcAccessOrder
2
,
// ABlockTransferSrcVectorDim
8
,
// ABlockTransferSrcScalarPerVector
8
,
// ABlockTransferDstScalarPerVector_AK1
1
,
// ABlockLdsExtraM
S
<
4
,
64
,
1
>
,
// BBlockTransferThreadClusterLengths_BK0_N_BK1
S
<
1
,
0
,
2
>
,
// BBlockTransferThreadClusterArrangeOrder
S
<
1
,
0
,
2
>
,
// BBlockTransferSrcAccessOrder
2
,
// BBlockTransferSrcVectorDim
8
,
// BBlockTransferSrcScalarPerVector
8
,
// BBlockTransferDstScalarPerVector_BK1
1
,
// BBlockLdsExtraN
1
,
1
,
S
<
1
,
32
,
1
,
8
>
,
8
>
;
const
ck
::
utils
::
conv
::
ConvParam
conv_param
{
3
,
1
,
16
,
16
,
8
,
{
3
,
3
,
3
},
{
17
,
17
,
17
},
{
2
,
2
,
2
},
{
1
,
1
,
1
},
{
1
,
1
,
1
},
{
1
,
1
,
1
}};
void
SetUp
()
override
{
if
(
!
ck
::
is_xdl_supported
())
{
GTEST_SKIP
();
}
}
template
<
typename
ADataType
,
typename
BDataType
>
bool
Run
(
ADataType
as
,
BDataType
bs
)
{
const
auto
in_g_n_c_wis_desc
=
ck
::
utils
::
conv
::
make_input_host_tensor_descriptor_g_n_c_wis_packed
<
InLayout
>
(
conv_param
);
const
auto
wei_g_k_c_xs_desc
=
ck
::
utils
::
conv
::
make_weight_host_tensor_descriptor_g_k_c_xs_packed
<
WeiLayout
>
(
conv_param
);
const
auto
out_g_n_k_wos_desc
=
ck
::
utils
::
conv
::
make_output_host_tensor_descriptor_g_n_k_wos_packed
<
OutLayout
>
(
conv_param
);
std
::
array
<
ck
::
index_t
,
NDimSpatial
+
3
>
a_g_n_c_wis_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
+
3
>
a_g_n_c_wis_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
+
3
>
b_g_k_c_xs_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
+
3
>
b_g_k_c_xs_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
+
3
>
e_g_n_k_wos_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
+
3
>
e_g_n_k_wos_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
conv_filter_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
conv_filter_dilations
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
input_left_pads
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
input_right_pads
{};
auto
copy
=
[](
const
auto
&
x
,
auto
&
y
)
{
ck
::
ranges
::
copy
(
x
,
y
.
begin
());
};
copy
(
in_g_n_c_wis_desc
.
GetLengths
(),
a_g_n_c_wis_lengths
);
copy
(
in_g_n_c_wis_desc
.
GetStrides
(),
a_g_n_c_wis_strides
);
copy
(
wei_g_k_c_xs_desc
.
GetLengths
(),
b_g_k_c_xs_lengths
);
copy
(
wei_g_k_c_xs_desc
.
GetStrides
(),
b_g_k_c_xs_strides
);
copy
(
out_g_n_k_wos_desc
.
GetLengths
(),
e_g_n_k_wos_lengths
);
copy
(
out_g_n_k_wos_desc
.
GetStrides
(),
e_g_n_k_wos_strides
);
copy
(
conv_param
.
conv_filter_strides_
,
conv_filter_strides
);
copy
(
conv_param
.
conv_filter_dilations_
,
conv_filter_dilations
);
copy
(
conv_param
.
input_left_pads_
,
input_left_pads
);
copy
(
conv_param
.
input_right_pads_
,
input_right_pads
);
std
::
array
<
const
void
*
,
0
>
ds
{};
// do Conv
auto
conv
=
DeviceGroupedConvNDMultiABFwdInstance
{};
auto
invoker
=
conv
.
MakeInvoker
();
auto
argument
=
conv
.
MakeArgument
(
as
,
bs
,
ds
,
nullptr
,
a_g_n_c_wis_lengths
,
a_g_n_c_wis_strides
,
b_g_k_c_xs_lengths
,
b_g_k_c_xs_strides
,
{},
{},
e_g_n_k_wos_lengths
,
e_g_n_k_wos_strides
,
conv_filter_strides
,
conv_filter_dilations
,
input_left_pads
,
input_right_pads
,
InElementOp
{},
WeiElementOp
{},
OutElementOp
{});
return
conv
.
IsSupportedArgument
(
argument
);
}
};
class
TestGroupedConvndFwdMultiAInterface
:
public
TestGroupedConvndFwdMultiABInterfaceBase
<
float
,
ck
::
Tuple
<
float
,
float
>
,
float
,
ScaleAdd
,
PassThrough
>
{
};
class
TestGroupedConvndFwdMultiBInterface
:
public
TestGroupedConvndFwdMultiABInterfaceBase
<
float
,
float
,
ck
::
Tuple
<
float
,
float
>
,
PassThrough
,
ScaleAdd
>
{
};
class
TestGroupedConvndFwdMultiABInterface
:
public
TestGroupedConvndFwdMultiABInterfaceBase
<
float
,
ck
::
Tuple
<
float
,
float
>
,
ck
::
Tuple
<
float
,
float
>
,
ScaleAdd
,
ScaleAdd
>
{
};
class
TestGroupedConvndFwdInterface
:
public
TestGroupedConvndFwdMultiABInterfaceBase
<
float
,
float
,
float
,
PassThrough
,
PassThrough
>
{
};
TEST_F
(
TestGroupedConvndFwdMultiAInterface
,
MultiA
)
{
std
::
array
<
const
void
*
,
NumAs
>
as
{
nullptr
,
nullptr
};
const
void
*
b
=
nullptr
;
EXPECT_TRUE
(
this
->
template
Run
(
as
,
b
));
}
TEST_F
(
TestGroupedConvndFwdMultiBInterface
,
MultiB
)
{
const
void
*
a
=
nullptr
;
std
::
array
<
const
void
*
,
NumBs
>
bs
{
nullptr
,
nullptr
};
EXPECT_TRUE
(
this
->
template
Run
(
a
,
bs
));
}
TEST_F
(
TestGroupedConvndFwdMultiABInterface
,
MultiAB
)
{
std
::
array
<
const
void
*
,
NumAs
>
as
{
nullptr
,
nullptr
};
std
::
array
<
const
void
*
,
NumBs
>
bs
{
nullptr
,
nullptr
};
EXPECT_TRUE
(
this
->
template
Run
(
as
,
bs
));
}
TEST_F
(
TestGroupedConvndFwdInterface
,
SingleAB
)
{
const
void
*
a
=
nullptr
;
const
void
*
b
=
nullptr
;
EXPECT_TRUE
(
this
->
template
Run
(
a
,
b
));
}
test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_d_interface_compatibility.cpp
0 → 100644
View file @
5903efe7
// SPDX-License-Identifier: MIT
// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iostream>
#include <initializer_list>
#include <tuple>
#include <vector>
#include "ck/ck.hpp"
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/tensor_operation/gpu/device/device_grouped_conv_fwd_multiple_d.hpp"
#include "ck/library/tensor_operation_instance/gpu/grouped_convolution_forward.hpp"
#include <gtest/gtest.h>
template
<
ck
::
index_t
...
Is
>
using
S
=
ck
::
Sequence
<
Is
...
>
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
class
TestGroupedConvndFwdMultiDInterfaceCompatibility
:
public
::
testing
::
Test
{
protected:
static
constexpr
ck
::
index_t
NDimSpatial
=
3
;
using
InDataType
=
float
;
using
WeiDataType
=
float
;
using
OutDataType
=
float
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
GNDHWC
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKZYXC
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
GNDHWK
;
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NDimSpatial
,
InLayout
,
WeiLayout
,
ck
::
Tuple
<>
,
OutLayout
,
InDataType
,
WeiDataType
,
ck
::
Tuple
<>
,
OutDataType
,
PassThrough
,
PassThrough
,
PassThrough
>
;
bool
Run
()
{
const
auto
op_ptrs
=
ck
::
tensor_operation
::
device
::
instance
::
DeviceOperationInstanceFactory
<
DeviceOp
>::
GetInstances
();
return
op_ptrs
.
size
()
!=
0
;
}
};
TEST_F
(
TestGroupedConvndFwdMultiDInterfaceCompatibility
,
CompatibilityTest
)
{
EXPECT_TRUE
(
this
->
Run
());
}
test/normalization/CMakeLists.txt
deleted
100644 → 0
View file @
2100ea4b
add_custom_target
(
test_normalization
)
add_gtest_executable
(
test_layernorm2d_fp32 test_layernorm2d_fp32.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_layernorm2d_fp32 PRIVATE utility device_normalization_instance
)
add_dependencies
(
test_normalization test_layernorm2d_fp32
)
endif
()
add_gtest_executable
(
test_groupnorm_fp32 test_groupnorm_fp32.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_groupnorm_fp32 PRIVATE utility device_normalization_instance
)
add_dependencies
(
test_normalization test_groupnorm_fp32
)
endif
()
add_gtest_executable
(
test_layernorm2d_fp16 test_layernorm2d_fp16.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_layernorm2d_fp16 PRIVATE utility device_normalization_instance
)
add_dependencies
(
test_normalization test_layernorm2d_fp16
)
endif
()
add_gtest_executable
(
test_groupnorm_fp16 test_groupnorm_fp16.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_groupnorm_fp16 PRIVATE utility device_normalization_instance
)
add_dependencies
(
test_normalization test_groupnorm_fp16
)
endif
()
test/normalization_fwd/CMakeLists.txt
0 → 100644
View file @
5903efe7
add_custom_target
(
test_normalization_fwd
)
add_gtest_executable
(
test_layernorm2d_fwd_fp32 test_layernorm2d_fwd_fp32.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_layernorm2d_fwd_fp32 PRIVATE utility device_normalization_fwd_instance
)
add_dependencies
(
test_normalization_fwd test_layernorm2d_fwd_fp32
)
endif
()
add_gtest_executable
(
test_groupnorm_fwd_fp32 test_groupnorm_fwd_fp32.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_groupnorm_fwd_fp32 PRIVATE utility device_normalization_fwd_instance
)
add_dependencies
(
test_normalization_fwd test_groupnorm_fwd_fp32
)
endif
()
add_gtest_executable
(
test_layernorm2d_fwd_fp16 test_layernorm2d_fwd_fp16.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_layernorm2d_fwd_fp16 PRIVATE utility device_normalization_fwd_instance
)
add_dependencies
(
test_normalization_fwd test_layernorm2d_fwd_fp16
)
endif
()
add_gtest_executable
(
test_layernorm4d_fwd_fp16 test_layernorm4d_fwd_fp16.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_layernorm4d_fwd_fp16 PRIVATE utility device_normalization_fwd_instance
)
add_dependencies
(
test_normalization_fwd test_layernorm4d_fwd_fp16
)
endif
()
add_gtest_executable
(
test_groupnorm_fwd_fp16 test_groupnorm_fwd_fp16.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_groupnorm_fwd_fp16 PRIVATE utility device_normalization_fwd_instance
)
add_dependencies
(
test_normalization_fwd test_groupnorm_fwd_fp16
)
endif
()
test/normalization/test_groupnorm_fp16.cpp
→
test/normalization
_fwd
/test_groupnorm_
fwd_
fp16.cpp
View file @
5903efe7
...
...
@@ -2,7 +2,7 @@
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_groupnorm_impl.hpp"
#include "profiler/profile_groupnorm_
fwd_
impl.hpp"
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
...
...
test/normalization/test_groupnorm_fp32.cpp
→
test/normalization
_fwd
/test_groupnorm_
fwd_
fp32.cpp
View file @
5903efe7
...
...
@@ -2,7 +2,7 @@
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_groupnorm_impl.hpp"
#include "profiler/profile_groupnorm_
fwd_
impl.hpp"
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
...
...
test/normalization/test_layernorm2d_fp16.cpp
→
test/normalization
_fwd
/test_layernorm2d_
fwd_
fp16.cpp
View file @
5903efe7
...
...
@@ -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 @
5903efe7
...
...
@@ -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 @
5903efe7
// 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
();
}
Prev
1
…
8
9
10
11
12
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