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
586d61ed
Commit
586d61ed
authored
May 17, 2023
by
rocking
Browse files
Modify the argument into vector.
parent
382c2ca0
Changes
6
Hide whitespace changes
Inline
Side-by-side
Showing
6 changed files
with
92 additions
and
85 deletions
+92
-85
example/13_pool2d_fwd/pool2d_fwd_common.hpp
example/13_pool2d_fwd/pool2d_fwd_common.hpp
+4
-4
example/48_pool3d_fwd/pool3d_fwd_common.hpp
example/48_pool3d_fwd/pool3d_fwd_common.hpp
+5
-6
include/ck/tensor_operation/gpu/device/device_pool_fwd.hpp
include/ck/tensor_operation/gpu/device/device_pool_fwd.hpp
+10
-10
include/ck/tensor_operation/gpu/device/impl/device_pool2d_fwd_nhwc_nhwc.hpp
...operation/gpu/device/impl/device_pool2d_fwd_nhwc_nhwc.hpp
+31
-27
include/ck/tensor_operation/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp
...eration/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp
+31
-27
library/include/ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp
...ary/reference_tensor_operation/cpu/reference_pool_fwd.hpp
+11
-11
No files found.
example/13_pool2d_fwd/pool2d_fwd_common.hpp
View file @
586d61ed
...
@@ -62,10 +62,10 @@ bool pool_test(bool do_verification,
...
@@ -62,10 +62,10 @@ bool pool_test(bool do_verification,
const
ck
::
index_t
Ho
=
(
Hi
+
in_left_pad_h
+
in_right_pad_h
-
Y
)
/
window_stride_h
+
1
;
const
ck
::
index_t
Ho
=
(
Hi
+
in_left_pad_h
+
in_right_pad_h
-
Y
)
/
window_stride_h
+
1
;
const
ck
::
index_t
Wo
=
(
Wi
+
in_left_pad_w
+
in_right_pad_w
-
X
)
/
window_stride_w
+
1
;
const
ck
::
index_t
Wo
=
(
Wi
+
in_left_pad_w
+
in_right_pad_w
-
X
)
/
window_stride_w
+
1
;
const
std
::
array
<
ck
::
index_t
,
2
>
window_spatial_lengths
{
{
Y
,
X
}
}
;
const
std
::
vector
<
ck
::
index_t
>
window_spatial_lengths
{
Y
,
X
};
const
std
::
array
<
ck
::
index_t
,
2
>
window_strides
{
{
window_stride_h
,
window_stride_w
}
}
;
const
std
::
vector
<
ck
::
index_t
>
window_strides
{
window_stride_h
,
window_stride_w
};
const
std
::
array
<
ck
::
index_t
,
2
>
input_left_pads
{
{
in_left_pad_h
,
in_left_pad_w
}
}
;
const
std
::
vector
<
ck
::
index_t
>
input_left_pads
{
in_left_pad_h
,
in_left_pad_w
};
const
std
::
array
<
ck
::
index_t
,
2
>
input_right_pads
{
{
in_right_pad_h
,
in_right_pad_w
}
}
;
const
std
::
vector
<
ck
::
index_t
>
input_right_pads
{
in_right_pad_h
,
in_right_pad_w
};
// tensor layout
// tensor layout
auto
f_host_tensor_descriptor
=
auto
f_host_tensor_descriptor
=
...
...
example/48_pool3d_fwd/pool3d_fwd_common.hpp
View file @
586d61ed
...
@@ -66,12 +66,11 @@ bool pool3d_test(bool do_verification,
...
@@ -66,12 +66,11 @@ bool pool3d_test(bool do_verification,
const
ck
::
index_t
Ho
=
(
Hi
+
in_left_pad_h
+
in_right_pad_h
-
Y
)
/
window_stride_h
+
1
;
const
ck
::
index_t
Ho
=
(
Hi
+
in_left_pad_h
+
in_right_pad_h
-
Y
)
/
window_stride_h
+
1
;
const
ck
::
index_t
Wo
=
(
Wi
+
in_left_pad_w
+
in_right_pad_w
-
X
)
/
window_stride_w
+
1
;
const
ck
::
index_t
Wo
=
(
Wi
+
in_left_pad_w
+
in_right_pad_w
-
X
)
/
window_stride_w
+
1
;
const
std
::
array
<
ck
::
index_t
,
3
>
window_spatial_lengths
{{
Z
,
Y
,
X
}};
const
std
::
vector
<
ck
::
index_t
>
window_spatial_lengths
{
Z
,
Y
,
X
};
const
std
::
array
<
ck
::
index_t
,
3
>
window_strides
{
const
std
::
vector
<
ck
::
index_t
>
window_strides
{
{
window_stride_d
,
window_stride_h
,
window_stride_w
}};
window_stride_d
,
window_stride_h
,
window_stride_w
};
const
std
::
array
<
ck
::
index_t
,
3
>
input_left_pads
{{
in_left_pad_d
,
in_left_pad_h
,
in_left_pad_w
}};
const
std
::
vector
<
ck
::
index_t
>
input_left_pads
{
in_left_pad_d
,
in_left_pad_h
,
in_left_pad_w
};
const
std
::
array
<
ck
::
index_t
,
3
>
input_right_pads
{
const
std
::
vector
<
ck
::
index_t
>
input_right_pads
{
in_right_pad_d
,
in_right_pad_h
,
in_right_pad_w
};
{
in_right_pad_d
,
in_right_pad_h
,
in_right_pad_w
}};
// tensor layout
// tensor layout
auto
f_host_tensor_descriptor
=
[](
std
::
size_t
N_
,
auto
f_host_tensor_descriptor
=
[](
std
::
size_t
N_
,
...
...
include/ck/tensor_operation/gpu/device/device_pool_fwd.hpp
View file @
586d61ed
...
@@ -26,16 +26,16 @@ struct DevicePoolFwd : public BaseOperator
...
@@ -26,16 +26,16 @@ struct DevicePoolFwd : public BaseOperator
MakeArgumentPointer
(
const
void
*
p_in_dev
,
MakeArgumentPointer
(
const
void
*
p_in_dev
,
void
*
p_out_dev
,
void
*
p_out_dev
,
void
*
p_out_indices_dev
,
void
*
p_out_indices_dev
,
std
::
array
<
ck
::
index_t
,
InOutRank
>
input_stride
,
std
::
vector
<
ck
::
index_t
>
input_stride
,
std
::
array
<
ck
::
index_t
,
InOutRank
>
output_stride
,
std
::
vector
<
ck
::
index_t
>
output_stride
,
std
::
array
<
ck
::
index_t
,
InOutRank
>
indices_stride
,
std
::
vector
<
ck
::
index_t
>
indices_stride
,
std
::
array
<
ck
::
index_t
,
InOutRank
>
input_lengths
,
std
::
vector
<
ck
::
index_t
>
input_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
window_lengths
,
std
::
vector
<
ck
::
index_t
>
window_lengths
,
std
::
array
<
ck
::
index_t
,
InOutRank
>
output_lengths
,
std
::
vector
<
ck
::
index_t
>
output_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
window_strides
,
std
::
vector
<
ck
::
index_t
>
window_strides
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_left_pads
,
std
::
vector
<
ck
::
index_t
>
input_left_pads
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_right_pads
,
std
::
vector
<
ck
::
index_t
>
input_right_pads
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
pooling_dims
)
=
0
;
std
::
vector
<
ck
::
index_t
>
pooling_dims
)
=
0
;
virtual
std
::
unique_ptr
<
BaseInvoker
>
MakeInvokerPointer
()
=
0
;
virtual
std
::
unique_ptr
<
BaseInvoker
>
MakeInvokerPointer
()
=
0
;
};
};
...
...
include/ck/tensor_operation/gpu/device/impl/device_pool2d_fwd_nhwc_nhwc.hpp
View file @
586d61ed
...
@@ -60,15 +60,14 @@ struct DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
...
@@ -60,15 +60,14 @@ struct DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
static
constexpr
ck
::
index_t
ReduceK_BlockTileSize
=
static
constexpr
ck
::
index_t
ReduceK_BlockTileSize
=
ReduceKThreadClusterSize
*
ReduceKThreadSliceSize
;
ReduceKThreadClusterSize
*
ReduceKThreadSliceSize
;
static
auto
static
auto
MakeABGridDescriptor_A_M_K_B_M
(
ck
::
index_t
N
,
MakeABGridDescriptor_A_M_K_B_M
(
ck
::
index_t
N
,
ck
::
index_t
C
,
ck
::
index_t
C
,
std
::
vector
<
ck
::
index_t
>
input_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>
window_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
window_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>
output_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
output_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>
window_strides
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
window_strides
,
std
::
vector
<
ck
::
index_t
>
input_left_pads
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_left_pads
,
std
::
vector
<
ck
::
index_t
>
input_right_pads
)
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_right_pads
)
{
{
const
index_t
Hi
=
input_spatial_lengths
[
0
];
const
index_t
Hi
=
input_spatial_lengths
[
0
];
const
index_t
Wi
=
input_spatial_lengths
[
1
];
const
index_t
Wi
=
input_spatial_lengths
[
1
];
...
@@ -159,12 +158,12 @@ struct DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
...
@@ -159,12 +158,12 @@ struct DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
IndexDataType
*
p_out_indices_dev
,
IndexDataType
*
p_out_indices_dev
,
ck
::
index_t
N
,
ck
::
index_t
N
,
ck
::
index_t
C
,
ck
::
index_t
C
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
input_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>&
input_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
window_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>&
window_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
output_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>&
output_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
window_strides
,
std
::
vector
<
ck
::
index_t
>&
window_strides
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
input_left_pads
,
std
::
vector
<
ck
::
index_t
>&
input_left_pads
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
input_right_pads
)
std
::
vector
<
ck
::
index_t
>&
input_right_pads
)
:
p_in_dev_
{
p_in_dev
},
:
p_in_dev_
{
p_in_dev
},
p_out_dev_
{
p_out_dev
},
p_out_dev_
{
p_out_dev
},
p_out_indices_dev_
{
p_out_indices_dev
},
p_out_indices_dev_
{
p_out_indices_dev
},
...
@@ -286,17 +285,22 @@ struct DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
...
@@ -286,17 +285,22 @@ struct DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
MakeArgumentPointer
(
const
void
*
p_in_dev
,
MakeArgumentPointer
(
const
void
*
p_in_dev
,
void
*
p_out_dev
,
void
*
p_out_dev
,
void
*
p_out_indices_dev
,
void
*
p_out_indices_dev
,
std
::
array
<
ck
::
index_t
,
InOutRank
>
,
// Suppose tensor layout = NHWC
std
::
vector
<
ck
::
index_t
>
,
// Suppose tensor layout = NHWC
std
::
array
<
ck
::
index_t
,
InOutRank
>
,
// Suppose tensor layout = NHWC
std
::
vector
<
ck
::
index_t
>
,
// Suppose tensor layout = NHWC
std
::
array
<
ck
::
index_t
,
InOutRank
>
,
// Suppose tensor layout = NHWC
std
::
vector
<
ck
::
index_t
>
,
// Suppose tensor layout = NHWC
std
::
array
<
ck
::
index_t
,
InOutRank
>
input_lengths
,
std
::
vector
<
ck
::
index_t
>
input_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
window_lengths
,
std
::
vector
<
ck
::
index_t
>
window_lengths
,
std
::
array
<
ck
::
index_t
,
InOutRank
>
output_lengths
,
std
::
vector
<
ck
::
index_t
>
output_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
window_strides
,
std
::
vector
<
ck
::
index_t
>
window_strides
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_left_pads
,
std
::
vector
<
ck
::
index_t
>
input_left_pads
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_right_pads
,
std
::
vector
<
ck
::
index_t
>
input_right_pads
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
)
override
std
::
vector
<
ck
::
index_t
>
)
override
{
{
if
(
input_lengths
.
size
()
!=
InOutRank
||
window_lengths
.
size
()
!=
WindowRank
||
input_lengths
.
size
()
!=
InOutRank
||
window_strides
.
size
()
!=
WindowRank
||
input_left_pads
.
size
()
!=
WindowRank
||
input_right_pads
.
size
()
!=
WindowRank
)
throw
std
::
runtime_error
(
"dimension is incorrect"
);
index_t
N
=
input_lengths
[
0
];
index_t
N
=
input_lengths
[
0
];
index_t
C
=
input_lengths
[
1
];
index_t
C
=
input_lengths
[
1
];
index_t
Hi
=
input_lengths
[
2
];
index_t
Hi
=
input_lengths
[
2
];
...
@@ -304,8 +308,8 @@ struct DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
...
@@ -304,8 +308,8 @@ struct DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
index_t
Ho
=
output_lengths
[
2
];
index_t
Ho
=
output_lengths
[
2
];
index_t
Wo
=
output_lengths
[
3
];
index_t
Wo
=
output_lengths
[
3
];
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_spatial_lengths
=
{
Hi
,
Wi
};
std
::
vector
<
ck
::
index_t
>
input_spatial_lengths
=
{
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
WindowRank
>
output_spatial_lengths
=
{
Ho
,
Wo
};
std
::
vector
<
ck
::
index_t
>
output_spatial_lengths
=
{
Ho
,
Wo
};
return
std
::
make_unique
<
Argument
>
(
static_cast
<
const
InDataType
*>
(
p_in_dev
),
return
std
::
make_unique
<
Argument
>
(
static_cast
<
const
InDataType
*>
(
p_in_dev
),
static_cast
<
OutDataType
*>
(
p_out_dev
),
static_cast
<
OutDataType
*>
(
p_out_dev
),
...
...
include/ck/tensor_operation/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp
View file @
586d61ed
...
@@ -58,15 +58,14 @@ struct DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
...
@@ -58,15 +58,14 @@ struct DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
static
constexpr
ck
::
index_t
M_BlockTileSize
=
MThreadClusterSize
*
MThreadSliceSize
;
static
constexpr
ck
::
index_t
M_BlockTileSize
=
MThreadClusterSize
*
MThreadSliceSize
;
static
constexpr
ck
::
index_t
K_BlockTileSize
=
KThreadClusterSize
*
KThreadSliceSize
;
static
constexpr
ck
::
index_t
K_BlockTileSize
=
KThreadClusterSize
*
KThreadSliceSize
;
static
auto
static
auto
MakeABGridDescriptor_A_M_K_B_M
(
ck
::
index_t
N
,
MakeABGridDescriptor_A_M_K_B_M
(
ck
::
index_t
N
,
ck
::
index_t
C
,
ck
::
index_t
C
,
std
::
vector
<
ck
::
index_t
>
input_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>
window_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
window_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>
output_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
output_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>
window_strides
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
window_strides
,
std
::
vector
<
ck
::
index_t
>
input_left_pads
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_left_pads
,
std
::
vector
<
ck
::
index_t
>
input_right_pads
)
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_right_pads
)
{
{
const
index_t
Di
=
input_spatial_lengths
[
0
];
const
index_t
Di
=
input_spatial_lengths
[
0
];
const
index_t
Hi
=
input_spatial_lengths
[
1
];
const
index_t
Hi
=
input_spatial_lengths
[
1
];
...
@@ -165,12 +164,12 @@ struct DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
...
@@ -165,12 +164,12 @@ struct DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
IndexDataType
*
p_out_indices_dev
,
IndexDataType
*
p_out_indices_dev
,
ck
::
index_t
N
,
ck
::
index_t
N
,
ck
::
index_t
C
,
ck
::
index_t
C
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
input_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>&
input_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
window_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>&
window_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
output_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>&
output_spatial_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
window_strides
,
std
::
vector
<
ck
::
index_t
>&
window_strides
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
input_left_pads
,
std
::
vector
<
ck
::
index_t
>&
input_left_pads
,
std
::
array
<
ck
::
index_t
,
WindowRank
>&
input_right_pads
)
std
::
vector
<
ck
::
index_t
>&
input_right_pads
)
:
p_in_dev_
{
p_in_dev
},
:
p_in_dev_
{
p_in_dev
},
p_out_dev_
{
p_out_dev
},
p_out_dev_
{
p_out_dev
},
p_out_indices_dev_
{
p_out_indices_dev
},
p_out_indices_dev_
{
p_out_indices_dev
},
...
@@ -291,17 +290,22 @@ struct DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
...
@@ -291,17 +290,22 @@ struct DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
MakeArgumentPointer
(
const
void
*
p_in_dev
,
MakeArgumentPointer
(
const
void
*
p_in_dev
,
void
*
p_out_dev
,
void
*
p_out_dev
,
void
*
p_out_indices_dev
,
void
*
p_out_indices_dev
,
std
::
array
<
ck
::
index_t
,
InOutRank
>
,
// Suppose tensor layout = NDHWC
std
::
vector
<
ck
::
index_t
>
,
// Suppose tensor layout = NDHWC
std
::
array
<
ck
::
index_t
,
InOutRank
>
,
// Suppose tensor layout = NDHWC
std
::
vector
<
ck
::
index_t
>
,
// Suppose tensor layout = NDHWC
std
::
array
<
ck
::
index_t
,
InOutRank
>
,
// Suppose tensor layout = NDHWC
std
::
vector
<
ck
::
index_t
>
,
// Suppose tensor layout = NDHWC
std
::
array
<
ck
::
index_t
,
InOutRank
>
input_lengths
,
std
::
vector
<
ck
::
index_t
>
input_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
window_lengths
,
std
::
vector
<
ck
::
index_t
>
window_lengths
,
std
::
array
<
ck
::
index_t
,
InOutRank
>
output_lengths
,
std
::
vector
<
ck
::
index_t
>
output_lengths
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
window_strides
,
std
::
vector
<
ck
::
index_t
>
window_strides
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_left_pads
,
std
::
vector
<
ck
::
index_t
>
input_left_pads
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_right_pads
,
std
::
vector
<
ck
::
index_t
>
input_right_pads
,
std
::
array
<
ck
::
index_t
,
WindowRank
>
)
override
std
::
vector
<
ck
::
index_t
>
)
override
{
{
if
(
input_lengths
.
size
()
!=
InOutRank
||
window_lengths
.
size
()
!=
WindowRank
||
input_lengths
.
size
()
!=
InOutRank
||
window_strides
.
size
()
!=
WindowRank
||
input_left_pads
.
size
()
!=
WindowRank
||
input_right_pads
.
size
()
!=
WindowRank
)
throw
std
::
runtime_error
(
"dimension is incorrect"
);
index_t
N
=
input_lengths
[
0
];
index_t
N
=
input_lengths
[
0
];
index_t
C
=
input_lengths
[
1
];
index_t
C
=
input_lengths
[
1
];
index_t
Di
=
input_lengths
[
2
];
index_t
Di
=
input_lengths
[
2
];
...
@@ -311,8 +315,8 @@ struct DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
...
@@ -311,8 +315,8 @@ struct DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C
index_t
Ho
=
output_lengths
[
3
];
index_t
Ho
=
output_lengths
[
3
];
index_t
Wo
=
output_lengths
[
4
];
index_t
Wo
=
output_lengths
[
4
];
std
::
array
<
ck
::
index_t
,
WindowRank
>
input_spatial_lengths
=
{
Di
,
Hi
,
Wi
};
std
::
vector
<
ck
::
index_t
>
input_spatial_lengths
=
{
Di
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
WindowRank
>
output_spatial_lengths
=
{
Do
,
Ho
,
Wo
};
std
::
vector
<
ck
::
index_t
>
output_spatial_lengths
=
{
Do
,
Ho
,
Wo
};
return
std
::
make_unique
<
Argument
>
(
static_cast
<
const
InDataType
*>
(
p_in_dev
),
return
std
::
make_unique
<
Argument
>
(
static_cast
<
const
InDataType
*>
(
p_in_dev
),
static_cast
<
OutDataType
*>
(
p_out_dev
),
static_cast
<
OutDataType
*>
(
p_out_dev
),
...
...
library/include/ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp
View file @
586d61ed
...
@@ -37,10 +37,10 @@ struct ReferencePoolingFwd : public device::BaseOperator
...
@@ -37,10 +37,10 @@ struct ReferencePoolingFwd : public device::BaseOperator
Argument
(
const
Tensor
<
InDataType
>&
in
,
Argument
(
const
Tensor
<
InDataType
>&
in
,
Tensor
<
OutDataType
>&
out
,
Tensor
<
OutDataType
>&
out
,
Tensor
<
IndexDataType
>&
out_indices
,
Tensor
<
IndexDataType
>&
out_indices
,
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
window_spatial_lengths
,
const
std
::
vector
<
ck
::
index_t
>&
window_spatial_lengths
,
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
window_strides
,
const
std
::
vector
<
ck
::
index_t
>&
window_strides
,
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
in_left_pads
,
const
std
::
vector
<
ck
::
index_t
>&
in_left_pads
,
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
/*in_right_pads*/
)
const
std
::
vector
<
ck
::
index_t
>&
/*in_right_pads*/
)
:
in_
(
in
),
:
in_
(
in
),
out_
(
out
),
out_
(
out
),
out_indices_
(
out_indices
),
out_indices_
(
out_indices
),
...
@@ -56,9 +56,9 @@ struct ReferencePoolingFwd : public device::BaseOperator
...
@@ -56,9 +56,9 @@ struct ReferencePoolingFwd : public device::BaseOperator
const
Tensor
<
InDataType
>&
in_
;
const
Tensor
<
InDataType
>&
in_
;
Tensor
<
OutDataType
>&
out_
;
Tensor
<
OutDataType
>&
out_
;
Tensor
<
IndexDataType
>&
out_indices_
;
Tensor
<
IndexDataType
>&
out_indices_
;
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
window_spatial_lengths_
;
const
std
::
vector
<
ck
::
index_t
>&
window_spatial_lengths_
;
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
window_strides_
;
const
std
::
vector
<
ck
::
index_t
>&
window_strides_
;
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
in_left_pads_
;
const
std
::
vector
<
ck
::
index_t
>&
in_left_pads_
;
int
reduceLength_
;
int
reduceLength_
;
};
};
...
@@ -306,10 +306,10 @@ struct ReferencePoolingFwd : public device::BaseOperator
...
@@ -306,10 +306,10 @@ struct ReferencePoolingFwd : public device::BaseOperator
static
auto
MakeArgument
(
const
Tensor
<
InDataType
>&
in
,
static
auto
MakeArgument
(
const
Tensor
<
InDataType
>&
in
,
Tensor
<
OutDataType
>&
out
,
Tensor
<
OutDataType
>&
out
,
Tensor
<
IndexDataType
>&
out_indices
,
Tensor
<
IndexDataType
>&
out_indices
,
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
window_spatial_lengths
,
const
std
::
vector
<
ck
::
index_t
>&
window_spatial_lengths
,
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
window_strides
,
const
std
::
vector
<
ck
::
index_t
>&
window_strides
,
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
in_left_pads
,
const
std
::
vector
<
ck
::
index_t
>&
in_left_pads
,
const
std
::
array
<
ck
::
index_t
,
WindowRank
>&
in_right_pads
)
const
std
::
vector
<
ck
::
index_t
>&
in_right_pads
)
{
{
return
Argument
{
in
,
return
Argument
{
in
,
out
,
out
,
...
...
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