Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
BaiXuePrincess
Paddle
提交
c52ed8de
P
Paddle
项目概览
BaiXuePrincess
/
Paddle
与 Fork 源项目一致
Fork自
PaddlePaddle / Paddle
通知
1
Star
1
Fork
0
代码
文件
提交
分支
Tags
贡献者
分支图
Diff
Issue
0
列表
看板
标记
里程碑
合并请求
0
Wiki
0
Wiki
分析
仓库
DevOps
项目成员
Pages
P
Paddle
项目概览
项目概览
详情
发布
仓库
仓库
文件
提交
分支
标签
贡献者
分支图
比较
Issue
0
Issue
0
列表
看板
标记
里程碑
合并请求
0
合并请求
0
Pages
分析
分析
仓库分析
DevOps
Wiki
0
Wiki
成员
成员
收起侧边栏
关闭侧边栏
动态
分支图
创建新Issue
提交
Issue看板
提交
c52ed8de
编写于
11月 29, 2017
作者:
S
sweetsky0901
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
format code
上级
bd561384
变更
6
隐藏空白更改
内联
并排
Showing
6 changed file
with
54 addition
and
82 deletion
+54
-82
paddle/operators/math/unpooling.cc
paddle/operators/math/unpooling.cc
+9
-13
paddle/operators/math/unpooling.cu
paddle/operators/math/unpooling.cu
+23
-36
paddle/operators/math/unpooling.h
paddle/operators/math/unpooling.h
+9
-14
paddle/operators/unpool_op.cc
paddle/operators/unpool_op.cc
+7
-8
paddle/operators/unpool_op.cu.cc
paddle/operators/unpool_op.cu.cc
+6
-8
paddle/operators/unpool_op.h
paddle/operators/unpool_op.h
+0
-3
未找到文件。
paddle/operators/math/unpooling.cc
浏览文件 @
c52ed8de
...
@@ -13,17 +13,15 @@ See the License for the specific language governing permissions and
...
@@ -13,17 +13,15 @@ See the License for the specific language governing permissions and
limitations under the License. */
limitations under the License. */
#include "paddle/operators/math/unpooling.h"
#include "paddle/operators/math/unpooling.h"
namespace
paddle
{
namespace
paddle
{
namespace
operators
{
namespace
operators
{
namespace
math
{
namespace
math
{
// All tensors are in NCHW format
template
<
typename
T
>
template
<
typename
T
>
class
Unpool2dMaxFunctor
<
platform
::
CPUPlace
,
T
>
{
class
Unpool2dMaxFunctor
<
platform
::
CPUPlace
,
T
>
{
public:
public:
void
operator
()(
const
platform
::
DeviceContext
&
context
,
void
operator
()(
const
framework
::
Tensor
&
input
,
const
platform
::
DeviceContext
&
context
,
const
framework
::
Tensor
&
input
,
const
framework
::
Tensor
&
indices
,
framework
::
Tensor
*
output
)
{
const
framework
::
Tensor
&
indices
,
framework
::
Tensor
*
output
)
{
const
int
batch_size
=
input
.
dims
()[
0
];
const
int
batch_size
=
input
.
dims
()[
0
];
const
int
input_height
=
input
.
dims
()[
2
];
const
int
input_height
=
input
.
dims
()[
2
];
const
int
input_width
=
input
.
dims
()[
3
];
const
int
input_width
=
input
.
dims
()[
3
];
...
@@ -51,13 +49,11 @@ public:
...
@@ -51,13 +49,11 @@ public:
};
};
template
<
class
T
>
template
<
class
T
>
class
Unpool2dMaxGradFunctor
<
platform
::
CPUPlace
,
T
>
{
class
Unpool2dMaxGradFunctor
<
platform
::
CPUPlace
,
T
>
{
public:
public:
void
operator
()(
const
platform
::
DeviceContext
&
context
,
void
operator
()(
const
framework
::
Tensor
&
input
,
const
platform
::
DeviceContext
&
context
,
const
framework
::
Tensor
&
input
,
const
framework
::
Tensor
&
indices
,
const
framework
::
Tensor
&
indices
,
const
framework
::
Tensor
&
output
,
const
framework
::
Tensor
&
output
,
const
framework
::
Tensor
&
output_grad
,
framework
::
Tensor
*
input_grad
)
{
const
framework
::
Tensor
&
output_grad
,
framework
::
Tensor
*
input_grad
)
{
const
int
batch_size
=
input
.
dims
()[
0
];
const
int
batch_size
=
input
.
dims
()[
0
];
const
int
input_height
=
input
.
dims
()[
2
];
const
int
input_height
=
input
.
dims
()[
2
];
const
int
input_width
=
input
.
dims
()[
3
];
const
int
input_width
=
input
.
dims
()[
3
];
...
...
paddle/operators/math/unpooling.cu
浏览文件 @
c52ed8de
...
@@ -19,14 +19,10 @@ namespace paddle {
...
@@ -19,14 +19,10 @@ namespace paddle {
namespace
operators
{
namespace
operators
{
namespace
math
{
namespace
math
{
template
<
typename
T
>
template
<
typename
T
>
__global__
void
KernelUnpool2dMax
(
const
int
nthreads
,
const
T
*
input_data
,
__global__
void
KernelUnpool2dMax
(
const
int
*
indices_data
,
const
int
nthreads
,
const
T
*
input_data
,
const
int
*
indices_data
,
const
int
input_height
,
const
int
input_height
,
const
int
input_width
,
const
int
channels
,
const
int
input_width
,
T
*
output_data
,
const
int
output_height
,
const
int
output_width
)
{
const
int
channels
,
T
*
output_data
,
const
int
output_height
,
const
int
output_width
)
{
int
in_n_stride
=
input_height
*
input_width
*
channels
;
int
in_n_stride
=
input_height
*
input_width
*
channels
;
int
in_c_stride
=
input_height
*
input_width
;
int
in_c_stride
=
input_height
*
input_width
;
int
out_n_stride
=
output_height
*
output_width
*
channels
;
int
out_n_stride
=
output_height
*
output_width
*
channels
;
...
@@ -44,16 +40,11 @@ __global__ void KernelUnpool2dMax(const int nthreads, const T* input_data,
...
@@ -44,16 +40,11 @@ __global__ void KernelUnpool2dMax(const int nthreads, const T* input_data,
}
}
}
}
template
<
typename
T
>
template
<
typename
T
>
__global__
void
KernelUnpool2dMaxGrad
(
const
int
nthreads
,
const
T
*
input_data
,
__global__
void
KernelUnpool2dMaxGrad
(
const
int
*
indices_data
,
const
int
nthreads
,
const
T
*
input_data
,
const
int
*
indices_data
,
const
int
input_height
,
const
int
input_height
,
const
int
input_width
,
const
int
channels
,
const
int
input_width
,
const
T
*
output_data
,
const
T
*
output_grad
,
const
int
output_height
,
const
int
channels
,
const
int
output_width
,
T
*
input_grad
)
{
const
T
*
output_data
,
const
T
*
output_grad
,
const
int
output_height
,
const
int
output_width
,
T
*
input_grad
)
{
int
in_n_stride
=
input_height
*
input_width
*
channels
;
int
in_n_stride
=
input_height
*
input_width
*
channels
;
int
in_c_stride
=
input_height
*
input_width
;
int
in_c_stride
=
input_height
*
input_width
;
int
out_n_stride
=
output_height
*
output_width
*
channels
;
int
out_n_stride
=
output_height
*
output_width
*
channels
;
...
@@ -75,11 +66,10 @@ __global__ void KernelUnpool2dMaxGrad(const int nthreads, const T* input_data,
...
@@ -75,11 +66,10 @@ __global__ void KernelUnpool2dMaxGrad(const int nthreads, const T* input_data,
*/
*/
template
<
typename
T
>
template
<
typename
T
>
class
Unpool2dMaxFunctor
<
platform
::
GPUPlace
,
T
>
{
class
Unpool2dMaxFunctor
<
platform
::
GPUPlace
,
T
>
{
public:
public:
void
operator
()(
const
platform
::
DeviceContext
&
context
,
void
operator
()(
const
framework
::
Tensor
&
input
,
const
platform
::
DeviceContext
&
context
,
const
framework
::
Tensor
&
input
,
const
framework
::
Tensor
&
indices
,
const
framework
::
Tensor
&
indices
,
framework
::
Tensor
*
output
)
{
framework
::
Tensor
*
output
)
{
const
int
batch_size
=
input
.
dims
()[
0
];
const
int
batch_size
=
input
.
dims
()[
0
];
const
int
input_height
=
input
.
dims
()[
2
];
const
int
input_height
=
input
.
dims
()[
2
];
const
int
input_width
=
input
.
dims
()[
3
];
const
int
input_width
=
input
.
dims
()[
3
];
...
@@ -91,12 +81,11 @@ public:
...
@@ -91,12 +81,11 @@ public:
T
*
output_data
=
output
->
mutable_data
<
T
>
(
context
.
GetPlace
());
T
*
output_data
=
output
->
mutable_data
<
T
>
(
context
.
GetPlace
());
int
threads
=
1024
;
int
threads
=
1024
;
int
grid
=
(
input
.
numel
()
+
threads
-
1
)
/
threads
;
int
grid
=
(
input
.
numel
()
+
threads
-
1
)
/
threads
;
KernelUnpool2dMax
<
KernelUnpool2dMax
<
T
><<<
grid
,
threads
,
0
,
T
><<<
grid
,
threads
,
0
,
reinterpret_cast
<
const
platform
::
CUDADeviceContext
&>
(
context
)
reinterpret_cast
<
const
platform
::
CUDADeviceContext
&>
(
context
)
.
stream
()
>>>
(
input
.
numel
(),
input_data
,
indices_data
,
.
stream
()
>>>
(
input
.
numel
(),
input_data
,
indices_data
,
input_height
,
input_width
,
output_channels
,
input_height
,
input_width
,
output_channels
,
output_data
,
output_height
,
output_width
);
output_data
,
output_height
,
output_width
);
}
}
};
};
/*
/*
...
@@ -104,7 +93,7 @@ public:
...
@@ -104,7 +93,7 @@ public:
*/
*/
template
<
typename
T
>
template
<
typename
T
>
class
Unpool2dMaxGradFunctor
<
platform
::
GPUPlace
,
T
>
{
class
Unpool2dMaxGradFunctor
<
platform
::
GPUPlace
,
T
>
{
public:
public:
void
operator
()(
const
platform
::
DeviceContext
&
context
,
void
operator
()(
const
platform
::
DeviceContext
&
context
,
const
framework
::
Tensor
&
input
,
const
framework
::
Tensor
&
input
,
const
framework
::
Tensor
&
indices
,
const
framework
::
Tensor
&
indices
,
...
@@ -124,13 +113,11 @@ public:
...
@@ -124,13 +113,11 @@ public:
T
*
input_grad_data
=
input_grad
->
mutable_data
<
T
>
(
context
.
GetPlace
());
T
*
input_grad_data
=
input_grad
->
mutable_data
<
T
>
(
context
.
GetPlace
());
int
threads
=
1024
;
int
threads
=
1024
;
int
grid
=
(
input
.
numel
()
+
threads
-
1
)
/
threads
;
int
grid
=
(
input
.
numel
()
+
threads
-
1
)
/
threads
;
KernelUnpool2dMaxGrad
<
KernelUnpool2dMaxGrad
<
T
><<<
grid
,
threads
,
0
,
T
><<<
grid
,
threads
,
0
,
reinterpret_cast
<
const
platform
::
CUDADeviceContext
&>
(
context
)
reinterpret_cast
<
const
platform
::
CUDADeviceContext
&>
(
context
)
.
stream
()
>>>
(
input
.
numel
(),
input_data
,
indices_data
,
.
stream
()
>>>
(
input
.
numel
(),
input_data
,
indices_data
,
input_height
,
input_width
,
output_channels
,
output_data
,
input_height
,
input_width
,
output_channels
,
output_grad_data
,
output_height
,
output_width
,
input_grad_data
);
output_data
,
output_grad_data
,
output_height
,
output_width
,
input_grad_data
);
}
}
};
};
template
class
Unpool2dMaxGradFunctor
<
platform
::
GPUPlace
,
float
>;
template
class
Unpool2dMaxGradFunctor
<
platform
::
GPUPlace
,
float
>;
...
...
paddle/operators/math/unpooling.h
浏览文件 @
c52ed8de
...
@@ -18,25 +18,20 @@ limitations under the License. */
...
@@ -18,25 +18,20 @@ limitations under the License. */
namespace
paddle
{
namespace
paddle
{
namespace
operators
{
namespace
operators
{
namespace
math
{
namespace
math
{
template
<
typename
Place
,
typename
T
>
template
<
typename
Place
,
typename
T
>
class
Unpool2dMaxFunctor
{
class
Unpool2dMaxFunctor
{
public:
public:
void
operator
()(
const
platform
::
DeviceContext
&
context
,
void
operator
()(
const
framework
::
Tensor
&
input
,
const
platform
::
DeviceContext
&
context
,
const
framework
::
Tensor
&
input
,
const
framework
::
Tensor
&
indices
,
framework
::
Tensor
*
output
);
const
framework
::
Tensor
&
indices
,
framework
::
Tensor
*
output
);
};
};
template
<
typename
Place
,
class
T
>
template
<
typename
Place
,
class
T
>
class
Unpool2dMaxGradFunctor
{
class
Unpool2dMaxGradFunctor
{
public:
public:
void
operator
()(
const
platform
::
DeviceContext
&
context
,
void
operator
()(
const
framework
::
Tensor
&
input
,
const
platform
::
DeviceContext
&
context
,
const
framework
::
Tensor
&
input
,
const
framework
::
Tensor
&
indices
,
const
framework
::
Tensor
&
indices
,
const
framework
::
Tensor
&
output
,
const
framework
::
Tensor
&
output
,
const
framework
::
Tensor
&
output_grad
,
framework
::
Tensor
*
input_grad
);
const
framework
::
Tensor
&
output_grad
,
framework
::
Tensor
*
input_grad
);
};
};
}
// namespace math
}
// namespace math
}
// namespace operators
}
// namespace operators
...
...
paddle/operators/unpool_op.cc
浏览文件 @
c52ed8de
...
@@ -31,13 +31,12 @@ class Unpool2dOpMaker : public framework::OpProtoAndCheckerMaker {
...
@@ -31,13 +31,12 @@ class Unpool2dOpMaker : public framework::OpProtoAndCheckerMaker {
"(Tensor) The input tensor of the indices given out by MaxPool2d. "
"(Tensor) The input tensor of the indices given out by MaxPool2d. "
"The format of input tensor is NCHW. Where N is batch size, C is the "
"The format of input tensor is NCHW. Where N is batch size, C is the "
"number of channels, H and W is the height and width of feature."
);
"number of channels, H and W is the height and width of feature."
);
AddOutput
(
AddOutput
(
"Out"
,
"Out"
,
"(Tensor) The output tensor of unpool operator."
"(Tensor) The output tensor of unpool operator."
"The format of output tensor is also NCHW."
"The format of output tensor is also NCHW."
"Where N is batch size, C is "
"Where N is batch size, C is "
"the number of channels, H and W is the height and "
"the number of channels, H and W is the height and "
"width of feature."
);
"width of feature."
);
AddAttr
<
std
::
vector
<
int
>>
(
AddAttr
<
std
::
vector
<
int
>>
(
"ksize"
,
"ksize"
,
"(vector), the unpooling window size(height, width) "
"(vector), the unpooling window size(height, width) "
...
@@ -138,7 +137,7 @@ namespace ops = paddle::operators;
...
@@ -138,7 +137,7 @@ namespace ops = paddle::operators;
REGISTER_OP
(
unpool
,
ops
::
UnpoolOp
,
ops
::
Unpool2dOpMaker
,
unpool_grad
,
REGISTER_OP
(
unpool
,
ops
::
UnpoolOp
,
ops
::
Unpool2dOpMaker
,
unpool_grad
,
ops
::
UnpoolOpGrad
);
ops
::
UnpoolOpGrad
);
REGISTER_OP_CPU_KERNEL
(
REGISTER_OP_CPU_KERNEL
(
unpool
,
ops
::
UnpoolKernel
<
paddle
::
platform
::
CPUPlace
,
float
>
,
unpool
,
ops
::
UnpoolKernel
<
paddle
::
platform
::
CPUPlace
,
float
>
,
ops
::
UnpoolKernel
<
paddle
::
platform
::
CPUPlace
,
double
>
);
ops
::
UnpoolKernel
<
paddle
::
platform
::
CPUPlace
,
double
>
);
REGISTER_OP_CPU_KERNEL
(
REGISTER_OP_CPU_KERNEL
(
unpool_grad
,
ops
::
UnpoolGradKernel
<
paddle
::
platform
::
CPUPlace
,
float
>
,
unpool_grad
,
ops
::
UnpoolGradKernel
<
paddle
::
platform
::
CPUPlace
,
float
>
,
...
...
paddle/operators/unpool_op.cu.cc
浏览文件 @
c52ed8de
...
@@ -15,11 +15,9 @@ limitations under the License. */
...
@@ -15,11 +15,9 @@ limitations under the License. */
#include "paddle/operators/unpool_op.h"
#include "paddle/operators/unpool_op.h"
namespace
ops
=
paddle
::
operators
;
namespace
ops
=
paddle
::
operators
;
REGISTER_OP_GPU_KERNEL
(
unpool
,
REGISTER_OP_GPU_KERNEL
(
ops
::
UnpoolKernel
<
paddle
::
platform
::
GPUPlace
,
float
>
,
unpool
,
ops
::
UnpoolKernel
<
paddle
::
platform
::
GPUPlace
,
float
>
,
ops
::
UnpoolKernel
<
paddle
::
platform
::
GPUPlace
,
double
>
);
ops
::
UnpoolKernel
<
paddle
::
platform
::
GPUPlace
,
double
>
);
REGISTER_OP_GPU_KERNEL
(
unpool_grad
,
REGISTER_OP_GPU_KERNEL
(
ops
::
UnpoolGradKernel
<
paddle
::
platform
::
GPUPlace
,
unpool_grad
,
ops
::
UnpoolGradKernel
<
paddle
::
platform
::
GPUPlace
,
float
>
,
float
>
,
ops
::
UnpoolGradKernel
<
paddle
::
platform
::
GPUPlace
,
double
>
);
ops
::
UnpoolGradKernel
<
paddle
::
platform
::
GPUPlace
,
double
>
);
paddle/operators/unpool_op.h
浏览文件 @
c52ed8de
...
@@ -20,7 +20,6 @@ limitations under the License. */
...
@@ -20,7 +20,6 @@ limitations under the License. */
namespace
paddle
{
namespace
paddle
{
namespace
operators
{
namespace
operators
{
template
<
typename
Place
,
typename
T
>
template
<
typename
Place
,
typename
T
>
class
UnpoolKernel
:
public
framework
::
OpKernel
<
T
>
{
class
UnpoolKernel
:
public
framework
::
OpKernel
<
T
>
{
public:
public:
...
@@ -41,7 +40,6 @@ class UnpoolKernel : public framework::OpKernel<T> {
...
@@ -41,7 +40,6 @@ class UnpoolKernel : public framework::OpKernel<T> {
unpool2d_max_forward
(
context
.
device_context
(),
*
in_x
,
*
in_y
,
out
);
unpool2d_max_forward
(
context
.
device_context
(),
*
in_x
,
*
in_y
,
out
);
}
}
};
};
template
<
typename
Place
,
typename
T
>
template
<
typename
Place
,
typename
T
>
class
UnpoolGradKernel
:
public
framework
::
OpKernel
<
T
>
{
class
UnpoolGradKernel
:
public
framework
::
OpKernel
<
T
>
{
public:
public:
...
@@ -69,6 +67,5 @@ class UnpoolGradKernel : public framework::OpKernel<T> {
...
@@ -69,6 +67,5 @@ class UnpoolGradKernel : public framework::OpKernel<T> {
*
out_grad
,
in_x_grad
);
*
out_grad
,
in_x_grad
);
}
}
};
};
}
// namespace operators
}
// namespace operators
}
// namespace paddle
}
// namespace paddle
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录