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
cf326690
Commit
cf326690
authored
Apr 20, 2022
by
rocking
Browse files
[What] Use F32 as the acc of reduce sum
[Why] Prevent loss of precision
parent
c16f789d
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
10 additions
and
10 deletions
+10
-10
example/19_gemm_softmax/gemm_softmax_xdl_fp16.cpp
example/19_gemm_softmax/gemm_softmax_xdl_fp16.cpp
+10
-10
No files found.
example/19_gemm_softmax/gemm_softmax_xdl_fp16.cpp
View file @
cf326690
...
...
@@ -98,15 +98,15 @@ constexpr ck::ReduceTensorOp ReduceMaxId = ck::ReduceTensorOp::MAX;
constexpr
ck
::
ReduceTensorOp
ReduceSumId
=
ck
::
ReduceTensorOp
::
ADD
;
constexpr
bool
ReducePropagateNan
=
false
;
using
ReduceMaxOp
=
typename
ck
::
reduce_binary_operator
<
CDataType
,
ReduceMaxId
>::
opType
;
using
ReduceSumOp
=
typename
ck
::
reduce_binary_operator
<
C
DataType
,
ReduceSumId
>::
opType
;
using
ReduceSumOp
=
typename
ck
::
reduce_binary_operator
<
Acc
DataType
,
ReduceSumId
>::
opType
;
using
ReduceMaxInElementwiseOperation
=
typename
ck
::
reduce_unary_operator
<
CDataType
,
ReduceMaxId
,
true
,
true
>::
InElementwiseOperation
;
using
ReduceMaxAccElementwiseOperation
=
typename
ck
::
reduce_unary_operator
<
CDataType
,
ReduceMaxId
,
true
,
true
>::
AccElementwiseOperation
;
using
ReduceSumInElementwiseOperation
=
typename
ck
::
reduce_unary_operator
<
C
DataType
,
ReduceSumId
,
true
,
true
>::
InElementwiseOperation
;
typename
ck
::
reduce_unary_operator
<
Acc
DataType
,
ReduceSumId
,
true
,
true
>::
InElementwiseOperation
;
using
ReduceSumAccElementwiseOperation
=
typename
ck
::
reduce_unary_operator
<
C
DataType
,
ReduceSumId
,
true
,
true
>::
AccElementwiseOperation
;
typename
ck
::
reduce_unary_operator
<
Acc
DataType
,
ReduceSumId
,
true
,
true
>::
AccElementwiseOperation
;
using
DeviceReduceMaxInstance
=
ck
::
tensor_operation
::
device
::
DeviceReduceBlockWise
<
CDataType
,
...
...
@@ -130,7 +130,7 @@ using DeviceReduceMaxInstance =
using
DeviceReduceSumInstance
=
ck
::
tensor_operation
::
device
::
DeviceReduceBlockWise
<
CDataType
,
C
DataType
,
Acc
DataType
,
CDataType
,
Rank
,
NumReduceDim
,
...
...
@@ -188,7 +188,7 @@ using HostReduceMaxInstance = ReductionHost<HostReduceDataType,
false
>
;
using
HostReduceSumInstance
=
ReductionHost
<
HostReduceDataType
,
HostReduce
DataType
,
Acc
DataType
,
HostReduceDataType
,
ReduceSumId
,
Rank
,
...
...
@@ -504,15 +504,15 @@ int main(int argc, char* argv[])
bool
result
=
true
;
if
(
result
&=
ck
::
utils
::
check_err
(
c_m_n
.
mData
,
host_c_m_n
.
mData
))
std
::
cout
<<
"[PASS] -
c_m_n
"
<<
std
::
endl
;
std
::
cout
<<
"[PASS] -
gemm
"
<<
std
::
endl
;
if
(
result
&=
ck
::
utils
::
check_err
(
c_n_max
.
mData
,
host_c_n_max
.
mData
))
std
::
cout
<<
"[PASS] -
c_n_
max"
<<
std
::
endl
;
std
::
cout
<<
"[PASS] -
reduce
max"
<<
std
::
endl
;
if
(
result
&=
ck
::
utils
::
check_err
(
exp_m_n
.
mData
,
host_exp_m_n
.
mData
))
std
::
cout
<<
"[PASS] -
exp_m_n
"
<<
std
::
endl
;
std
::
cout
<<
"[PASS] -
broadcast sub + exp
"
<<
std
::
endl
;
if
(
result
&=
ck
::
utils
::
check_err
(
exp_n_sum
.
mData
,
host_exp_n_sum
.
mData
))
std
::
cout
<<
"[PASS] -
exp_n_
sum"
<<
std
::
endl
;
std
::
cout
<<
"[PASS] -
reduce
sum"
<<
std
::
endl
;
if
(
result
&=
ck
::
utils
::
check_err
(
softmax_m_n
.
mData
,
host_softmax_m_n
.
mData
))
std
::
cout
<<
"[PASS] -
softmax_m_n
"
<<
std
::
endl
;
std
::
cout
<<
"[PASS] -
broadcast div
"
<<
std
::
endl
;
}
return
0
;
}
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