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
e1914e7f
Commit
e1914e7f
authored
May 18, 2023
by
rocking
Browse files
Rename AccDatatype to ComputeDatatype
parent
c89fb586
Changes
11
Hide whitespace changes
Inline
Side-by-side
Showing
11 changed files
with
49 additions
and
49 deletions
+49
-49
example/13_pool2d_fwd/pool2d_fwd_common.hpp
example/13_pool2d_fwd/pool2d_fwd_common.hpp
+3
-3
example/13_pool2d_fwd/pool2d_fwd_fp16.cpp
example/13_pool2d_fwd/pool2d_fwd_fp16.cpp
+2
-2
example/13_pool2d_fwd/pool2d_fwd_fp32.cpp
example/13_pool2d_fwd/pool2d_fwd_fp32.cpp
+2
-2
example/48_pool3d_fwd/pool3d_fwd_common.hpp
example/48_pool3d_fwd/pool3d_fwd_common.hpp
+3
-3
example/48_pool3d_fwd/pool3d_fwd_fp16.cpp
example/48_pool3d_fwd/pool3d_fwd_fp16.cpp
+2
-2
include/ck/tensor_operation/gpu/device/impl/device_pool2d_fwd_nhwc_nhwc.hpp
...operation/gpu/device/impl/device_pool2d_fwd_nhwc_nhwc.hpp
+3
-3
include/ck/tensor_operation/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp
...eration/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp
+3
-3
library/include/ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp
...ary/reference_tensor_operation/cpu/reference_pool_fwd.hpp
+19
-19
library/src/tensor_operation_instance/gpu/pool_fwd/pool_fwd_instance_common.hpp
...ration_instance/gpu/pool_fwd/pool_fwd_instance_common.hpp
+8
-8
profiler/include/profiler/profile_pool2d_fwd_impl.hpp
profiler/include/profiler/profile_pool2d_fwd_impl.hpp
+2
-2
profiler/include/profiler/profile_pool3d_fwd_impl.hpp
profiler/include/profiler/profile_pool3d_fwd_impl.hpp
+2
-2
No files found.
example/13_pool2d_fwd/pool2d_fwd_common.hpp
View file @
e1914e7f
...
...
@@ -21,7 +21,7 @@
template
<
typename
InDataType
,
typename
OutDataType
,
typename
Acc
DataType
,
typename
Compute
DataType
,
typename
IndexDataType
,
typename
InLayout
,
typename
OutLayout
,
...
...
@@ -49,7 +49,7 @@ bool pool_test(bool do_verification,
InDataType
,
// InDataType
OutDataType
,
// OutDataType
IndexDataType
,
// IndexDataType
Acc
DataType
,
//
Acc
DataType
Compute
DataType
,
//
Compute
DataType
ReduceOpId
,
OutputIndex
,
64
,
// BlockSize
...
...
@@ -156,7 +156,7 @@ bool pool_test(bool do_verification,
2
,
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
ReduceOpId
,
PropagateNan
,
...
...
example/13_pool2d_fwd/pool2d_fwd_fp16.cpp
View file @
e1914e7f
...
...
@@ -11,7 +11,7 @@
using
InDataType
=
ck
::
half_t
;
using
OutDataType
=
ck
::
half_t
;
using
Acc
DataType
=
float
;
using
Compute
DataType
=
float
;
using
IndexDataType
=
int32_t
;
...
...
@@ -90,7 +90,7 @@ int main(int argc, char* argv[])
bool
pass
=
pool_test
<
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
InLayout
,
OutLayout
,
...
...
example/13_pool2d_fwd/pool2d_fwd_fp32.cpp
View file @
e1914e7f
...
...
@@ -11,7 +11,7 @@
using
InDataType
=
float
;
using
OutDataType
=
float
;
using
Acc
DataType
=
float
;
using
Compute
DataType
=
float
;
using
IndexDataType
=
int32_t
;
...
...
@@ -90,7 +90,7 @@ int main(int argc, char* argv[])
bool
pass
=
pool_test
<
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
InLayout
,
OutLayout
,
...
...
example/48_pool3d_fwd/pool3d_fwd_common.hpp
View file @
e1914e7f
...
...
@@ -20,7 +20,7 @@
template
<
typename
InDataType
,
typename
OutDataType
,
typename
Acc
DataType
,
typename
Compute
DataType
,
typename
IndexDataType
,
typename
InLayout
,
typename
OutLayout
,
...
...
@@ -52,7 +52,7 @@ bool pool3d_test(bool do_verification,
InDataType
,
// InDataType
OutDataType
,
// OutDataType
IndexDataType
,
// IndexDataType
Acc
DataType
,
//
Acc
DataType
Compute
DataType
,
//
Compute
DataType
ReduceOpId
,
OutputIndex
,
64
,
// BlockSize
...
...
@@ -152,7 +152,7 @@ bool pool3d_test(bool do_verification,
3
,
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
ReduceOpId
,
PropagateNan
,
...
...
example/48_pool3d_fwd/pool3d_fwd_fp16.cpp
View file @
e1914e7f
...
...
@@ -11,7 +11,7 @@
using
InDataType
=
ck
::
half_t
;
using
OutDataType
=
ck
::
half_t
;
using
Acc
DataType
=
float
;
using
Compute
DataType
=
float
;
using
IndexDataType
=
int32_t
;
...
...
@@ -53,7 +53,7 @@ int main()
bool
pass
=
pool3d_test
<
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
InLayout
,
OutLayout
,
...
...
include/ck/tensor_operation/gpu/device/impl/device_pool2d_fwd_nhwc_nhwc.hpp
View file @
e1914e7f
...
...
@@ -21,7 +21,7 @@ namespace device {
template
<
typename
InDataType
,
typename
OutDataType
,
typename
IndexDataType
,
// enable if OutputIndex == true
typename
Acc
DataType
,
typename
Compute
DataType
,
ck
::
ReduceTensorOp
ReduceOpId
,
bool
OutputIndex
,
ck
::
index_t
BlockSize
,
...
...
@@ -211,7 +211,7 @@ struct DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
using
gridwise_reduce
=
GridwiseReduction_mk_to_m_threadwise
<
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
AGridDesc_M_K
,
BGridDesc_M
,
...
...
@@ -234,7 +234,7 @@ struct DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
false
,
// don't have index input
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
AGridDesc_M_K
,
BGridDesc_M
,
...
...
include/ck/tensor_operation/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp
View file @
e1914e7f
...
...
@@ -21,7 +21,7 @@ namespace device {
template
<
typename
InDataType
,
typename
OutDataType
,
typename
IndexDataType
,
// enable if OutputIndex == true
typename
Acc
DataType
,
typename
Compute
DataType
,
ck
::
ReduceTensorOp
ReduceOpId
,
bool
OutputIndex
,
ck
::
index_t
BlockSize
,
...
...
@@ -216,7 +216,7 @@ struct DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
using
gridwise_reduce
=
GridwiseReduction_mk_to_m_threadwise
<
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
AGridDesc_M_K
,
BGridDesc_M
,
...
...
@@ -239,7 +239,7 @@ struct DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
false
,
// don't have index input
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
AGridDesc_M_K
,
BGridDesc_M
,
...
...
library/include/ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp
View file @
e1914e7f
...
...
@@ -22,7 +22,7 @@ template <index_t InOutRank,
index_t
WindowRank
,
typename
InDataType
,
typename
OutDataType
,
typename
Acc
DataType
,
typename
Compute
DataType
,
typename
IndexDataType
,
ck
::
ReduceTensorOp
ReduceOpId
,
bool
PropagateNan
,
...
...
@@ -77,11 +77,11 @@ struct ReferencePoolingFwd : public device::BaseOperator
if
constexpr
(
!
OutputIndex
)
{
using
Accumulation
=
ck
::
detail
::
AccumulateWithNanCheck
<
PropagateNan
,
ReduceOperation
,
Acc
DataType
>
;
using
Accumulation
=
ck
::
detail
::
AccumulateWithNanCheck
<
PropagateNan
,
ReduceOperation
,
Compute
DataType
>
;
auto
f_ncdhw
=
[
&
](
auto
n
,
auto
c
,
auto
do_
,
auto
ho
,
auto
wo
)
{
auto
accuVal
=
ReduceOperation
::
template
GetIdentityValue
<
Acc
DataType
>();
auto
accuVal
=
ReduceOperation
::
template
GetIdentityValue
<
Compute
DataType
>();
for
(
ck
::
index_t
z
=
0
;
z
<
arg
.
window_spatial_lengths_
[
0
];
++
z
)
{
...
...
@@ -100,8 +100,8 @@ struct ReferencePoolingFwd : public device::BaseOperator
wi
>=
0
&&
wi
<
static_cast
<
ck
::
index_t
>
(
arg
.
in_
.
mDesc
.
GetLengths
()[
4
]))
{
Acc
DataType
currVal
=
static_cast
<
Acc
DataType
>
(
arg
.
in_
(
n
,
c
,
di
,
hi
,
wi
));
Compute
DataType
currVal
=
static_cast
<
Compute
DataType
>
(
arg
.
in_
(
n
,
c
,
di
,
hi
,
wi
));
in_elementwise_op
(
currVal
,
currVal
);
...
...
@@ -127,11 +127,11 @@ struct ReferencePoolingFwd : public device::BaseOperator
{
using
Accumulation
=
ck
::
detail
::
AccumulateWithIndexAndNanCheck
<
PropagateNan
,
ReduceOperation
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
>
;
auto
f_ncdhw
=
[
&
](
auto
n
,
auto
c
,
auto
do_
,
auto
ho
,
auto
wo
)
{
auto
accuVal
=
ReduceOperation
::
template
GetIdentityValue
<
Acc
DataType
>();
auto
accuVal
=
ReduceOperation
::
template
GetIdentityValue
<
Compute
DataType
>();
IndexDataType
accuIndex
=
0
;
for
(
ck
::
index_t
z
=
0
;
z
<
arg
.
window_spatial_lengths_
[
0
];
++
z
)
...
...
@@ -151,8 +151,8 @@ struct ReferencePoolingFwd : public device::BaseOperator
wi
>=
0
&&
wi
<
static_cast
<
ck
::
index_t
>
(
arg
.
in_
.
mDesc
.
GetLengths
()[
4
]))
{
Acc
DataType
currVal
=
static_cast
<
Acc
DataType
>
(
arg
.
in_
(
n
,
c
,
di
,
hi
,
wi
));
Compute
DataType
currVal
=
static_cast
<
Compute
DataType
>
(
arg
.
in_
(
n
,
c
,
di
,
hi
,
wi
));
IndexDataType
currIndex
=
arg
.
in_
.
GetOffsetFromMultiIndex
(
n
,
c
,
di
,
hi
,
wi
);
...
...
@@ -194,11 +194,11 @@ struct ReferencePoolingFwd : public device::BaseOperator
if
constexpr
(
!
OutputIndex
)
{
using
Accumulation
=
ck
::
detail
::
AccumulateWithNanCheck
<
PropagateNan
,
ReduceOperation
,
Acc
DataType
>
;
using
Accumulation
=
ck
::
detail
::
AccumulateWithNanCheck
<
PropagateNan
,
ReduceOperation
,
Compute
DataType
>
;
auto
f_nchw
=
[
&
](
auto
n
,
auto
c
,
auto
ho
,
auto
wo
)
{
auto
accuVal
=
ReduceOperation
::
template
GetIdentityValue
<
Acc
DataType
>();
auto
accuVal
=
ReduceOperation
::
template
GetIdentityValue
<
Compute
DataType
>();
for
(
ck
::
index_t
y
=
0
;
y
<
arg
.
window_spatial_lengths_
[
0
];
++
y
)
{
...
...
@@ -211,8 +211,8 @@ struct ReferencePoolingFwd : public device::BaseOperator
wi
>=
0
&&
wi
<
static_cast
<
ck
::
index_t
>
(
arg
.
in_
.
mDesc
.
GetLengths
()[
3
]))
{
Acc
DataType
currVal
=
static_cast
<
Acc
DataType
>
(
arg
.
in_
(
n
,
c
,
hi
,
wi
));
Compute
DataType
currVal
=
static_cast
<
Compute
DataType
>
(
arg
.
in_
(
n
,
c
,
hi
,
wi
));
in_elementwise_op
(
currVal
,
currVal
);
...
...
@@ -236,11 +236,11 @@ struct ReferencePoolingFwd : public device::BaseOperator
{
using
Accumulation
=
ck
::
detail
::
AccumulateWithIndexAndNanCheck
<
PropagateNan
,
ReduceOperation
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
>
;
auto
f_nchw
=
[
&
](
auto
n
,
auto
c
,
auto
ho
,
auto
wo
)
{
auto
accuVal
=
ReduceOperation
::
template
GetIdentityValue
<
Acc
DataType
>();
auto
accuVal
=
ReduceOperation
::
template
GetIdentityValue
<
Compute
DataType
>();
IndexDataType
accuIndex
=
0
;
for
(
ck
::
index_t
y
=
0
;
y
<
arg
.
window_spatial_lengths_
[
0
];
++
y
)
...
...
@@ -254,8 +254,8 @@ struct ReferencePoolingFwd : public device::BaseOperator
wi
>=
0
&&
wi
<
static_cast
<
ck
::
index_t
>
(
arg
.
in_
.
mDesc
.
GetLengths
()[
3
]))
{
Acc
DataType
currVal
=
static_cast
<
Acc
DataType
>
(
arg
.
in_
(
n
,
c
,
hi
,
wi
));
Compute
DataType
currVal
=
static_cast
<
Compute
DataType
>
(
arg
.
in_
(
n
,
c
,
hi
,
wi
));
IndexDataType
currIndex
=
arg
.
in_
.
GetOffsetFromMultiIndex
(
n
,
c
,
hi
,
wi
);
...
...
library/src/tensor_operation_instance/gpu/pool_fwd/pool_fwd_instance_common.hpp
View file @
e1914e7f
...
...
@@ -22,30 +22,30 @@ using F32 = float;
template
<
typename
InDataType
,
typename
OutDataType
,
typename
IndexDataType
,
typename
Acc
DataType
,
typename
Compute
DataType
,
ReduceTensorOp
ReduceOpId
,
bool
OutputIndex
>
using
device_pool2d_fwd_nhwc_instances
=
// clang-format off
std
::
tuple
<
DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Acc
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
1
,
1
,
1
>
,
DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Acc
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
2
,
1
,
2
>
,
DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Acc
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
4
,
1
,
4
>
DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Compute
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
1
,
1
,
1
>
,
DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Compute
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
2
,
1
,
2
>
,
DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Compute
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
4
,
1
,
4
>
// clang-format on
>
;
template
<
typename
InDataType
,
typename
OutDataType
,
typename
IndexDataType
,
typename
Acc
DataType
,
typename
Compute
DataType
,
ReduceTensorOp
ReduceOpId
,
bool
OutputIndex
>
using
device_pool3d_fwd_ndhwc_instances
=
// clang-format off
std
::
tuple
<
DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Acc
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
1
,
1
,
1
>
,
DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Acc
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
2
,
1
,
2
>
,
DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Acc
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
4
,
1
,
4
>
DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Compute
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
1
,
1
,
1
>
,
DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Compute
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
2
,
1
,
2
>
,
DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
<
InDataType
,
OutDataType
,
IndexDataType
,
Compute
DataType
,
ReduceOpId
,
OutputIndex
,
256
,
256
,
1
,
4
,
1
,
4
>
// clang-format on
>
;
...
...
profiler/include/profiler/profile_pool2d_fwd_impl.hpp
View file @
e1914e7f
...
...
@@ -19,7 +19,7 @@ namespace profiler {
template
<
typename
InDataType
,
typename
OutDataType
,
typename
Acc
DataType
,
typename
Compute
DataType
,
typename
IndexDataType
,
ck
::
ReduceTensorOp
ReduceOpId
,
bool
PropagateNan
,
...
...
@@ -119,7 +119,7 @@ bool profile_pool2d_fwd_impl(int do_verification,
WindowRank
,
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
ReduceOpId
,
PropagateNan
,
...
...
profiler/include/profiler/profile_pool3d_fwd_impl.hpp
View file @
e1914e7f
...
...
@@ -19,7 +19,7 @@ namespace profiler {
template
<
typename
InDataType
,
typename
OutDataType
,
typename
Acc
DataType
,
typename
Compute
DataType
,
typename
IndexDataType
,
ck
::
ReduceTensorOp
ReduceOpId
,
bool
PropagateNan
,
...
...
@@ -124,7 +124,7 @@ bool profile_pool3d_fwd_impl(int do_verification,
WindowRank
,
InDataType
,
OutDataType
,
Acc
DataType
,
Compute
DataType
,
IndexDataType
,
ReduceOpId
,
PropagateNan
,
...
...
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