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
36802715
Commit
36802715
authored
Apr 25, 2022
by
qinletao
Browse files
Merge branch 'develop' into add_mfma_f64
parents
ea970541
7c0b1498
Changes
44
Hide whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
16 additions
and
86 deletions
+16
-86
test/reduce/reduce_no_index.cpp
test/reduce/reduce_no_index.cpp
+4
-25
test/reduce/reduce_with_index.cpp
test/reduce/reduce_with_index.cpp
+5
-25
test/reference_conv_fwd/CMakeLists.txt
test/reference_conv_fwd/CMakeLists.txt
+1
-1
test/reference_conv_fwd/reference_conv_fwd.cpp
test/reference_conv_fwd/reference_conv_fwd.cpp
+6
-35
No files found.
test/reduce/reduce_no_index.cpp
View file @
36802715
...
...
@@ -37,19 +37,6 @@ static inline std::vector<int> get_invariant_dims(const std::vector<int>& reduce
return
invariantDims
;
};
// map the data type used by the GPU kernels to the corresponding type used by the host codes
template
<
typename
InType
>
struct
type_mapping
{
using
OutType
=
InType
;
};
template
<
>
struct
type_mapping
<
ck
::
half_t
>
{
using
OutType
=
half_float
::
half
;
};
constexpr
int
Rank
=
4
;
constexpr
ReduceTensorOp
ReduceOpId
=
ReduceTensorOp
::
AVG
;
...
...
@@ -226,13 +213,9 @@ bool test_reduce_no_index_impl(int init_method,
bool
result
=
true
;
using
HostInDataType
=
typename
type_mapping
<
InDataType
>::
OutType
;
using
HostOutDataType
=
typename
type_mapping
<
OutDataType
>::
OutType
;
using
HostAccDataType
=
typename
type_mapping
<
AccDataType
>::
OutType
;
ReductionHost
<
HostInDataType
,
HostAccDataType
,
HostOutDataType
,
ReductionHost
<
InDataType
,
AccDataType
,
OutDataType
,
ReduceOpId
,
Rank
,
NumReduceDim
,
...
...
@@ -240,11 +223,7 @@ bool test_reduce_no_index_impl(int init_method,
NeedIndices
>
hostReduce
(
in
.
mDesc
,
out_ref
.
mDesc
,
invariantDims
,
reduceDims
);
hostReduce
.
Run
(
alpha
,
reinterpret_cast
<
const
HostInDataType
*>
(
in
.
mData
.
data
()),
beta
,
reinterpret_cast
<
HostOutDataType
*>
(
out_ref
.
mData
.
data
()),
nullptr
);
hostReduce
.
Run
(
alpha
,
in
.
mData
.
data
(),
beta
,
out_ref
.
mData
.
data
(),
nullptr
);
const
auto
i_inLengths
=
to_int_vector
(
inLengths
);
const
auto
i_inStrides
=
to_int_vector
(
inStrides
);
...
...
test/reduce/reduce_with_index.cpp
View file @
36802715
...
...
@@ -36,19 +36,6 @@ static inline std::vector<int> get_invariant_dims(const std::vector<int>& reduce
return
invariantDims
;
};
// map the data type used by the GPU kernels to the corresponding type used by the host codes
template
<
typename
InType
>
struct
type_mapping
{
using
OutType
=
InType
;
};
template
<
>
struct
type_mapping
<
ck
::
half_t
>
{
using
OutType
=
half_float
::
half
;
};
constexpr
int
Rank
=
4
;
constexpr
ReduceTensorOp
ReduceOpId
=
ReduceTensorOp
::
AMAX
;
...
...
@@ -209,13 +196,9 @@ bool test_reduce_with_index_impl(int init_method,
bool
result
=
true
;
using
HostInDataType
=
typename
type_mapping
<
InDataType
>::
OutType
;
using
HostOutDataType
=
typename
type_mapping
<
OutDataType
>::
OutType
;
using
HostAccDataType
=
typename
type_mapping
<
AccDataType
>::
OutType
;
ReductionHost
<
HostInDataType
,
HostAccDataType
,
HostOutDataType
,
ReductionHost
<
InDataType
,
AccDataType
,
OutDataType
,
ReduceOpId
,
Rank
,
NumReduceDim
,
...
...
@@ -223,11 +206,8 @@ bool test_reduce_with_index_impl(int init_method,
NeedIndices
>
hostReduce
(
in
.
mDesc
,
out_ref
.
mDesc
,
invariantDims
,
reduceDims
);
hostReduce
.
Run
(
alpha
,
reinterpret_cast
<
const
HostInDataType
*>
(
in
.
mData
.
data
()),
beta
,
reinterpret_cast
<
HostOutDataType
*>
(
out_ref
.
mData
.
data
()),
out_indices_ref
.
mData
.
data
());
hostReduce
.
Run
(
alpha
,
in
.
mData
.
data
(),
beta
,
out_ref
.
mData
.
data
(),
out_indices_ref
.
mData
.
data
());
const
auto
i_inLengths
=
to_int_vector
(
inLengths
);
const
auto
i_inStrides
=
to_int_vector
(
inStrides
);
...
...
test/reference_conv_fwd/CMakeLists.txt
View file @
36802715
add_test_executable
(
test_reference_conv_fwd reference_conv_fwd.cpp
)
target_link_libraries
(
test_reference_conv_fwd PRIVATE host_tensor
)
target_link_libraries
(
test_reference_conv_fwd PRIVATE host_tensor
conv_fwd_util
)
test/reference_conv_fwd/reference_conv_fwd.cpp
View file @
36802715
#include <algorithm>
#include <cmath>
#include <cstdlib>
#include <half.hpp>
...
...
@@ -10,6 +9,7 @@
#include "config.hpp"
#include "conv_fwd_util.hpp"
#include "element_wise_operation.hpp"
#include "fill.hpp"
#include "host_tensor.hpp"
#include "reference_conv_fwd.hpp"
#include "tensor_layout.hpp"
...
...
@@ -19,35 +19,6 @@ using InElementOp = ck::tensor_operation::element_wise::PassThrough;
using
WeiElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
template
<
typename
T
>
struct
FillMonotonicSeq
{
T
m_init_value
{
0
};
T
m_step
{
1
};
template
<
typename
ForwardIter
>
void
operator
()(
ForwardIter
first
,
ForwardIter
last
)
const
{
std
::
generate
(
first
,
last
,
[
=
,
n
=
m_init_value
]()
mutable
{
auto
tmp
=
n
;
n
+=
m_step
;
return
tmp
;
});
}
};
template
<
typename
T
>
struct
FillConstant
{
T
m_value
{
0
};
template
<
typename
ForwardIter
>
void
operator
()(
ForwardIter
first
,
ForwardIter
last
)
const
{
std
::
fill
(
first
,
last
,
m_value
);
}
};
template
<
ck
::
index_t
NDim
,
typename
InDataType
=
float
,
typename
WeiDataType
=
float
,
...
...
@@ -55,8 +26,8 @@ template <ck::index_t NDim,
typename
InLayout
=
ck
::
tensor_layout
::
convolution
::
NHWC
,
typename
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
KYXC
,
typename
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NHWK
,
typename
FillInputOp
=
FillMonotonicSeq
<
InDataType
>,
typename
FillWeightsOp
=
FillConstant
<
WeiDataType
>>
typename
FillInputOp
=
ck
::
utils
::
FillMonotonicSeq
<
InDataType
>,
typename
FillWeightsOp
=
ck
::
utils
::
FillConstant
<
WeiDataType
>>
Tensor
<
OutDataType
>
run_reference_convolution_forward
(
const
ck
::
utils
::
conv
::
ConvParams
&
params
,
const
FillInputOp
&
fill_input_op
=
FillInputOp
{},
...
...
@@ -251,7 +222,7 @@ bool test_conv1d_nwc()
ck
::
tensor_layout
::
convolution
::
NWC
,
ck
::
tensor_layout
::
convolution
::
KXC
,
ck
::
tensor_layout
::
convolution
::
NWK
>
(
params
,
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
params
,
ck
::
utils
::
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
ref_dims
=
std
::
vector
<
std
::
size_t
>
{
2
,
16
,
16
};
ref_data
=
std
::
vector
<
float
>
{
...
...
@@ -349,7 +320,7 @@ bool test_conv3d_ncdhw()
ck
::
tensor_layout
::
convolution
::
NCDHW
,
ck
::
tensor_layout
::
convolution
::
KCZYX
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
(
params
,
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
params
,
ck
::
utils
::
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
std
::
vector
<
std
::
size_t
>
ref_dims
{
1
,
1
,
4
,
4
,
4
};
std
::
vector
<
float
>
ref_data
{
407.7
,
410.40002
,
413.09998
,
415.80002
,
423.90002
,
426.6
,
429.30002
,
432.
,
...
...
@@ -383,7 +354,7 @@ bool test_conv3d_ncdhw()
ck
::
tensor_layout
::
convolution
::
NCDHW
,
ck
::
tensor_layout
::
convolution
::
KCZYX
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
(
params
,
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
params
,
ck
::
utils
::
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
ref_dims
=
std
::
vector
<
std
::
size_t
>
{
1
,
2
,
4
,
4
,
4
};
ref_data
=
std
::
vector
<
float
>
{
2756.7002
,
2764.7998
,
2772.9001
,
2781.
,
2853.9001
,
2862.
,
2870.1
,
2878.2002
,
...
...
Prev
1
2
3
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