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
6b94f903
Commit
6b94f903
authored
Jul 24, 2023
by
rocking
Browse files
In to client example, this commit revise following:
1. Add dilation. 2. Use pool3d to implement pool2d
parent
a0f3b013
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
63 additions
and
15 deletions
+63
-15
client_example/19_pool_fwd/avg_pool3d_fwd.cpp
client_example/19_pool_fwd/avg_pool3d_fwd.cpp
+7
-3
client_example/19_pool_fwd/max_pool2d_fwd.cpp
client_example/19_pool_fwd/max_pool2d_fwd.cpp
+56
-12
No files found.
client_example/19_pool_fwd/avg_pool3d_fwd.cpp
View file @
6b94f903
...
@@ -65,9 +65,12 @@ int main(int argc, char* argv[])
...
@@ -65,9 +65,12 @@ int main(int argc, char* argv[])
ck
::
index_t
in_right_pad_h
=
1
;
ck
::
index_t
in_right_pad_h
=
1
;
ck
::
index_t
in_right_pad_w
=
1
;
ck
::
index_t
in_right_pad_w
=
1
;
ck
::
index_t
Do
=
(
Di
+
in_left_pad_d
+
in_right_pad_d
-
Z
)
/
window_stride_d
+
1
;
const
ck
::
index_t
Zs
=
(
Z
-
1
)
*
window_dilation_d
+
1
;
ck
::
index_t
Ho
=
(
Hi
+
in_left_pad_h
+
in_right_pad_h
-
Y
)
/
window_stride_h
+
1
;
const
ck
::
index_t
Ys
=
(
Y
-
1
)
*
window_dilation_h
+
1
;
ck
::
index_t
Wo
=
(
Wi
+
in_left_pad_w
+
in_right_pad_w
-
X
)
/
window_stride_w
+
1
;
const
ck
::
index_t
Xs
=
(
X
-
1
)
*
window_dilation_w
+
1
;
ck
::
index_t
Do
=
(
Di
+
in_left_pad_d
+
in_right_pad_d
-
Zs
)
/
window_stride_d
+
1
;
ck
::
index_t
Ho
=
(
Hi
+
in_left_pad_h
+
in_right_pad_h
-
Ys
)
/
window_stride_h
+
1
;
ck
::
index_t
Wo
=
(
Wi
+
in_left_pad_w
+
in_right_pad_w
-
Xs
)
/
window_stride_w
+
1
;
// Pool API only support the order of NCDHW
// Pool API only support the order of NCDHW
std
::
vector
<
ck
::
index_t
>
in_length
=
{
N
,
C
,
Di
,
Hi
,
Wi
};
std
::
vector
<
ck
::
index_t
>
in_length
=
{
N
,
C
,
Di
,
Hi
,
Wi
};
...
@@ -187,6 +190,7 @@ int main(int argc, char* argv[])
...
@@ -187,6 +190,7 @@ int main(int argc, char* argv[])
out_tensor_stride
,
out_tensor_stride
,
out_tensor_stride
,
out_tensor_stride
,
window_strides
,
window_strides
,
window_dilations
,
input_left_pads
,
input_left_pads
,
input_right_pads
,
input_right_pads
,
{
2
,
3
,
4
});
{
2
,
3
,
4
});
...
...
client_example/19_pool_fwd/max_pool2d_fwd.cpp
View file @
6b94f903
...
@@ -10,14 +10,15 @@
...
@@ -10,14 +10,15 @@
#include "ck/tensor_operation/gpu/device/device_pool_fwd.hpp"
#include "ck/tensor_operation/gpu/device/device_pool_fwd.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/library/tensor_operation_instance/gpu/pool
2
d_fwd.hpp"
#include "ck/library/tensor_operation_instance/gpu/pool
3
d_fwd.hpp"
using
InDataType
=
ck
::
half_t
;
using
InDataType
=
ck
::
half_t
;
using
OutDataType
=
ck
::
half_t
;
using
OutDataType
=
ck
::
half_t
;
using
IndexDataType
=
int32_t
;
using
IndexDataType
=
int32_t
;
constexpr
ck
::
index_t
InOutRank
=
4
;
// We use pool3d to implement pool2d in this example
constexpr
ck
::
index_t
WindowRank
=
2
;
constexpr
ck
::
index_t
InOutRank
=
5
;
constexpr
ck
::
index_t
WindowRank
=
3
;
#if 1
#if 1
constexpr
auto
ReduceOpId
=
ck
::
ReduceTensorOp
::
MAX
;
constexpr
auto
ReduceOpId
=
ck
::
ReduceTensorOp
::
MAX
;
constexpr
bool
OutputIndex
=
true
;
constexpr
bool
OutputIndex
=
true
;
...
@@ -42,6 +43,35 @@ struct SimpleDeviceMem
...
@@ -42,6 +43,35 @@ struct SimpleDeviceMem
void
*
p_mem_
;
void
*
p_mem_
;
};
};
void
TransformPool2dparamToPool3d
(
std
::
vector
<
ck
::
index_t
>&
input_lengths
,
std
::
vector
<
ck
::
index_t
>&
window_lengths
,
std
::
vector
<
ck
::
index_t
>&
output_lengths
,
std
::
vector
<
ck
::
index_t
>&
input_stride
,
std
::
vector
<
ck
::
index_t
>&
output_stride
,
std
::
vector
<
ck
::
index_t
>&
indices_stride
,
std
::
vector
<
ck
::
index_t
>&
window_strides
,
std
::
vector
<
ck
::
index_t
>&
window_dilations
,
std
::
vector
<
ck
::
index_t
>&
input_left_pads
,
std
::
vector
<
ck
::
index_t
>&
input_right_pads
,
std
::
vector
<
ck
::
index_t
>&
pooling_dims
)
{
// NCHW to NCDHW
input_lengths
.
insert
(
input_lengths
.
begin
()
+
2
,
1
);
output_lengths
.
insert
(
output_lengths
.
begin
()
+
2
,
1
);
input_stride
.
insert
(
input_stride
.
begin
()
+
2
,
0
);
output_stride
.
insert
(
output_stride
.
begin
()
+
2
,
0
);
indices_stride
.
insert
(
indices_stride
.
begin
()
+
2
,
0
);
// YX to ZYX
window_lengths
.
insert
(
window_lengths
.
begin
(),
1
);
window_strides
.
insert
(
window_strides
.
begin
(),
0
);
window_dilations
.
insert
(
window_dilations
.
begin
(),
0
);
input_left_pads
.
insert
(
input_left_pads
.
begin
(),
0
);
input_right_pads
.
insert
(
input_right_pads
.
begin
(),
0
);
pooling_dims
=
{
2
,
3
,
4
};
}
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
{
{
ck
::
index_t
N
=
2
;
ck
::
index_t
N
=
2
;
...
@@ -52,7 +82,6 @@ int main(int argc, char* argv[])
...
@@ -52,7 +82,6 @@ int main(int argc, char* argv[])
ck
::
index_t
Wi
=
30
;
ck
::
index_t
Wi
=
30
;
ck
::
index_t
window_stride_h
=
2
;
ck
::
index_t
window_stride_h
=
2
;
ck
::
index_t
window_stride_w
=
2
;
ck
::
index_t
window_stride_w
=
2
;
ck
::
index_t
window_dilation_d
=
1
;
ck
::
index_t
window_dilation_h
=
1
;
ck
::
index_t
window_dilation_h
=
1
;
ck
::
index_t
window_dilation_w
=
1
;
ck
::
index_t
window_dilation_w
=
1
;
ck
::
index_t
in_left_pad_h
=
1
;
ck
::
index_t
in_left_pad_h
=
1
;
...
@@ -60,18 +89,20 @@ int main(int argc, char* argv[])
...
@@ -60,18 +89,20 @@ int main(int argc, char* argv[])
ck
::
index_t
in_right_pad_h
=
1
;
ck
::
index_t
in_right_pad_h
=
1
;
ck
::
index_t
in_right_pad_w
=
1
;
ck
::
index_t
in_right_pad_w
=
1
;
ck
::
index_t
Ho
=
(
Hi
+
in_left_pad_h
+
in_right_pad_h
-
Y
)
/
window_stride_h
+
1
;
const
ck
::
index_t
Ys
=
(
Y
-
1
)
*
window_dilation_h
+
1
;
ck
::
index_t
Wo
=
(
Wi
+
in_left_pad_w
+
in_right_pad_w
-
X
)
/
window_stride_w
+
1
;
const
ck
::
index_t
Xs
=
(
X
-
1
)
*
window_dilation_w
+
1
;
ck
::
index_t
Ho
=
(
Hi
+
in_left_pad_h
+
in_right_pad_h
-
Ys
)
/
window_stride_h
+
1
;
ck
::
index_t
Wo
=
(
Wi
+
in_left_pad_w
+
in_right_pad_w
-
Xs
)
/
window_stride_w
+
1
;
// Pool API only support the order of NCHW
// Pool API only support the order of NCHW
std
::
vector
<
ck
::
index_t
>
in_length
=
{
N
,
C
,
Hi
,
Wi
};
std
::
vector
<
ck
::
index_t
>
in_length
=
{
N
,
C
,
Hi
,
Wi
};
std
::
vector
<
ck
::
index_t
>
out_length
=
{
N
,
C
,
Ho
,
Wo
};
std
::
vector
<
ck
::
index_t
>
out_length
=
{
N
,
C
,
Ho
,
Wo
};
std
::
vector
<
ck
::
index_t
>
window_spatial_lengths
=
{
Y
,
X
};
std
::
vector
<
ck
::
index_t
>
window_spatial_lengths
=
{
Y
,
X
};
std
::
vector
<
ck
::
index_t
>
window_strides
=
{
window_stride_h
,
window_stride_w
};
std
::
vector
<
ck
::
index_t
>
window_strides
=
{
window_stride_h
,
window_stride_w
};
std
::
vector
<
ck
::
index_t
>
window_dilations
{
std
::
vector
<
ck
::
index_t
>
window_dilations
=
{
window_dilation_h
,
window_dilation_w
};
window_dilation_d
,
window_dilation
_h
,
w
in
dow_dilation
_w
};
std
::
vector
<
ck
::
index_t
>
input_left_pads
=
{
in_left_pad
_h
,
in
_left_pad
_w
};
std
::
vector
<
ck
::
index_t
>
input_
lef
t_pads
=
{
in_
lef
t_pad_h
,
in_
lef
t_pad_w
};
std
::
vector
<
ck
::
index_t
>
input_
righ
t_pads
=
{
in_
righ
t_pad_h
,
in_
righ
t_pad_w
};
std
::
vector
<
ck
::
index_t
>
input_right_pads
=
{
in_right_pad_h
,
in_right_pad_w
};
std
::
vector
<
ck
::
index_t
>
pooling_dims
=
{
2
,
3
};
std
::
size_t
in_tensor_size
=
N
*
C
*
Hi
*
Wi
;
std
::
size_t
in_tensor_size
=
N
*
C
*
Hi
*
Wi
;
std
::
size_t
out_tensor_size
=
N
*
C
*
Ho
*
Wo
;
std
::
size_t
out_tensor_size
=
N
*
C
*
Ho
*
Wo
;
...
@@ -80,6 +111,18 @@ int main(int argc, char* argv[])
...
@@ -80,6 +111,18 @@ int main(int argc, char* argv[])
std
::
vector
<
ck
::
index_t
>
in_tensor_stride
=
{
C
*
Hi
*
Wi
,
1
,
Wi
*
C
,
C
};
std
::
vector
<
ck
::
index_t
>
in_tensor_stride
=
{
C
*
Hi
*
Wi
,
1
,
Wi
*
C
,
C
};
std
::
vector
<
ck
::
index_t
>
out_tensor_stride
=
{
C
*
Ho
*
Wo
,
1
,
Wo
*
C
,
C
};
std
::
vector
<
ck
::
index_t
>
out_tensor_stride
=
{
C
*
Ho
*
Wo
,
1
,
Wo
*
C
,
C
};
TransformPool2dparamToPool3d
(
in_length
,
window_spatial_lengths
,
out_length
,
in_tensor_stride
,
out_tensor_stride
,
out_tensor_stride
,
window_strides
,
window_dilations
,
input_left_pads
,
input_right_pads
,
pooling_dims
);
SimpleDeviceMem
in_device_buf
(
sizeof
(
InDataType
)
*
in_tensor_size
);
SimpleDeviceMem
in_device_buf
(
sizeof
(
InDataType
)
*
in_tensor_size
);
SimpleDeviceMem
out_device_buf
(
sizeof
(
OutDataType
)
*
out_tensor_size
);
SimpleDeviceMem
out_device_buf
(
sizeof
(
OutDataType
)
*
out_tensor_size
);
SimpleDeviceMem
out_indices_device_buf
(
sizeof
(
IndexDataType
)
*
out_tensor_size
);
SimpleDeviceMem
out_indices_device_buf
(
sizeof
(
IndexDataType
)
*
out_tensor_size
);
...
@@ -124,7 +167,7 @@ int main(int argc, char* argv[])
...
@@ -124,7 +167,7 @@ int main(int argc, char* argv[])
window_dilations
,
window_dilations
,
input_left_pads
,
input_left_pads
,
input_right_pads
,
input_right_pads
,
{
2
,
3
}
);
pooling_dims
);
auto
invoker_ptr
=
op_ptr
->
MakeInvokerPointer
();
auto
invoker_ptr
=
op_ptr
->
MakeInvokerPointer
();
...
@@ -181,9 +224,10 @@ int main(int argc, char* argv[])
...
@@ -181,9 +224,10 @@ int main(int argc, char* argv[])
out_tensor_stride
,
out_tensor_stride
,
out_tensor_stride
,
out_tensor_stride
,
window_strides
,
window_strides
,
window_dilations
,
input_left_pads
,
input_left_pads
,
input_right_pads
,
input_right_pads
,
{
2
,
3
}
);
pooling_dims
);
auto
invoker_ptr
=
op_ptr
->
MakeInvokerPointer
();
auto
invoker_ptr
=
op_ptr
->
MakeInvokerPointer
();
...
...
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