Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Oneflow-Inc
oneflow
提交
a717eb38
O
oneflow
项目概览
Oneflow-Inc
/
oneflow
上一次同步 接近 3 年
通知
13
Star
2733
Fork
0
代码
文件
提交
分支
Tags
贡献者
分支图
Diff
Issue
0
列表
看板
标记
里程碑
合并请求
0
DevOps
流水线
流水线任务
计划
Wiki
0
Wiki
分析
仓库
DevOps
项目成员
Pages
O
oneflow
项目概览
项目概览
详情
发布
仓库
仓库
文件
提交
分支
标签
贡献者
分支图
比较
Issue
0
Issue
0
列表
看板
标记
里程碑
合并请求
0
合并请求
0
Pages
DevOps
DevOps
流水线
流水线任务
计划
分析
分析
仓库分析
DevOps
Wiki
0
Wiki
成员
成员
收起侧边栏
关闭侧边栏
动态
分支图
创建新Issue
流水线任务
提交
Issue看板
体验新版 GitCode,发现更多精彩内容 >>
提交
a717eb38
编写于
7月 17, 2017
作者:
W
willzhang4a58
浏览文件
操作
浏览文件
下载
差异文件
add Model Update Op into OpMgr
上级
fdfdb39f
75453ebd
变更
5
隐藏空白更改
内联
并排
Showing
5 changed file
with
99 addition
and
7 deletion
+99
-7
oneflow/core/kernel/convolution_kernel.cu
oneflow/core/kernel/convolution_kernel.cu
+4
-2
oneflow/core/kernel/kernel_test_common.cpp
oneflow/core/kernel/kernel_test_common.cpp
+1
-1
oneflow/core/kernel/rmsprop_model_update_kernel.cpp
oneflow/core/kernel/rmsprop_model_update_kernel.cpp
+1
-2
oneflow/core/kernel/rmsprop_model_update_kernel.cu
oneflow/core/kernel/rmsprop_model_update_kernel.cu
+2
-2
oneflow/core/kernel/rmsprop_model_update_kernel_test.cpp
oneflow/core/kernel/rmsprop_model_update_kernel_test.cpp
+91
-0
未找到文件。
oneflow/core/kernel/convolution_kernel.cu
浏览文件 @
a717eb38
...
...
@@ -102,7 +102,8 @@ class ConvolutionKernelUtil<DeviceType::kGPU, FloatingPointType> final {
(
width
+
2
*
pad_w
-
(
dilation_w
*
(
kernel_w
-
1
)
+
1
))
/
stride_w
+
1
;
int
num_kernels
=
channels
*
height_col
*
width_col
;
Im2ColGpuKernel
<
FloatingPointType
>
<<<
BlocksNum4ThreadsNum
(
num_kernels
),
kCudaThreadsNumPerBlock
>>>
(
<<<
BlocksNum4ThreadsNum
(
num_kernels
),
kCudaThreadsNumPerBlock
,
0
,
ctx
.
device_ctx
->
cuda_stream
()
>>>
(
num_kernels
,
data_im
,
height
,
width
,
kernel_h
,
kernel_w
,
pad_h
,
pad_w
,
stride_h
,
stride_w
,
dilation_h
,
dilation_w
,
height_col
,
width_col
,
data_col
);
...
...
@@ -122,7 +123,8 @@ class ConvolutionKernelUtil<DeviceType::kGPU, FloatingPointType> final {
// To avoid involving atomic operations, we will launch one kernel per
// bottom dimension, and then in the kernel add up the top dimensions.
Col2ImGpuKernel
<
FloatingPointType
>
<<<
BlocksNum4ThreadsNum
(
num_kernels
),
kCudaThreadsNumPerBlock
>>>
(
<<<
BlocksNum4ThreadsNum
(
num_kernels
),
kCudaThreadsNumPerBlock
,
0
,
ctx
.
device_ctx
->
cuda_stream
()
>>>
(
num_kernels
,
data_col
,
height
,
width
,
channels
,
kernel_h
,
kernel_w
,
pad_h
,
pad_w
,
stride_h
,
stride_w
,
dilation_h
,
dilation_w
,
height_col
,
width_col
,
data_im
);
...
...
oneflow/core/kernel/kernel_test_common.cpp
浏览文件 @
a717eb38
...
...
@@ -62,7 +62,7 @@ class KernelTestCommon<DeviceType::kCPU, FloatingPointType> final {
size_t
dptr_size
=
lhs
->
shape
().
elem_cnt
();
for
(
size_t
i
=
0
;
i
<
dptr_size
;
++
i
)
{
ASSERT_
NEAR
(
dptr_lhs
[
i
],
dptr_rhs
[
i
],
0.0000001
);
ASSERT_
FLOAT_EQ
(
dptr_lhs
[
i
],
dptr_rhs
[
i
]
);
}
}
...
...
oneflow/core/kernel/rmsprop_model_update_kernel.cpp
浏览文件 @
a717eb38
...
...
@@ -56,8 +56,7 @@ class RMSPropMdUpdateKernelUtil<DeviceType::kCPU, FloatingPointType> final {
const
FloatingPointType
alpha
)
{
ctx
.
device_ctx
->
cpu_stream
()
->
SendWork
([
=
]()
{
for
(
int64_t
i
=
0
;
i
<
n
;
++
i
)
{
model
[
i
]
-=
alpha
*
model_diff
[
i
]
/
(
std
::
sqrt
(
mean_square
[
i
])
+
epsilon
);
model
[
i
]
-=
alpha
*
model_diff
[
i
]
/
std
::
sqrt
(
mean_square
[
i
]
+
epsilon
);
}
});
}
...
...
oneflow/core/kernel/rmsprop_model_update_kernel.cu
浏览文件 @
a717eb38
...
...
@@ -17,7 +17,7 @@ __global__ void UpdateMeanSquareGpu(const int64_t n,
}
}
// model -= alpha * model_diff /
(sqrt(mean_square)
+ epsilon)
// model -= alpha * model_diff /
sqrt(mean_square
+ epsilon)
template
<
typename
FloatingPointType
>
__global__
void
UpdateModelGpu
(
const
int64_t
n
,
FloatingPointType
*
model
,
const
FloatingPointType
*
model_diff
,
...
...
@@ -25,7 +25,7 @@ __global__ void UpdateModelGpu(const int64_t n, FloatingPointType* model,
const
FloatingPointType
epsilon
,
const
FloatingPointType
alpha
)
{
CUDA_1D_KERNEL_LOOP
(
i
,
n
)
{
model
[
i
]
-=
alpha
*
model_diff
[
i
]
/
(
std
::
sqrt
(
mean_square
[
i
])
+
epsilon
);
model
[
i
]
-=
alpha
*
model_diff
[
i
]
/
std
::
sqrt
(
mean_square
[
i
]
+
epsilon
);
}
}
...
...
oneflow/core/kernel/rmsprop_model_update_kernel_test.cpp
0 → 100644
浏览文件 @
a717eb38
#include "oneflow/core/kernel/rmsprop_model_update_kernel.h"
#include "oneflow/core/kernel/kernel_test_common.h"
namespace
oneflow
{
namespace
test
{
namespace
{
template
<
DeviceType
device_type
,
typename
FloatingPointType
>
Kernel
*
BuildRMSPropMdUpdateKernel
(
float
learning_rate
,
float
decay_rate
,
float
epsilon
)
{
OperatorConf
op_conf
;
op_conf
.
set_name
(
"rmsprop_model_update_test"
);
RMSPropModelUpdateOpConf
*
rmsprop_md_update_conf
=
op_conf
.
mutable_rmsprop_mdupdt_conf
();
rmsprop_md_update_conf
->
set_learning_rate
(
learning_rate
);
rmsprop_md_update_conf
->
set_decay_rate
(
decay_rate
);
rmsprop_md_update_conf
->
set_epsilon
(
epsilon
);
auto
rmsprop_md_update_op
=
ConstructOp
(
op_conf
);
OperatorProto
op_proto
;
rmsprop_md_update_op
->
ToProto
(
&
op_proto
);
auto
rmsprop_md_update_kernel
=
new
RMSPropMdUpdateKernel
<
device_type
,
FloatingPointType
>
();
rmsprop_md_update_kernel
->
InitFromOpProto
(
op_proto
);
return
rmsprop_md_update_kernel
;
}
void
InitJobDesc
(
int32_t
piece_size
,
int32_t
num_of_pieces_in_batch
)
{
JobConf
job_conf
;
job_conf
.
set_piece_size
(
piece_size
);
job_conf
.
set_num_of_pieces_in_batch
(
num_of_pieces_in_batch
);
JobDesc
::
Singleton
()
->
InitFromJobConf
(
job_conf
);
}
template
<
DeviceType
device_type
,
typename
FloatingPointType
>
std
::
function
<
Blob
*
(
const
std
::
string
&
)
>
BuildBnInOp2BlobPtr
()
{
using
KTCommon
=
KernelTestCommon
<
device_type
,
FloatingPointType
>
;
std
::
vector
<
int64_t
>
dim_vec
=
{
1
,
3
,
2
};
auto
bn2blob_ptr
=
new
HashMap
<
std
::
string
,
Blob
*>
;
(
*
bn2blob_ptr
)[
"model"
]
=
KTCommon
::
CreateBlobWithSameValue
(
dim_vec
,
3
);
(
*
bn2blob_ptr
)[
"mean_square"
]
=
KTCommon
::
CreateBlobWithSameValue
(
dim_vec
,
2
);
(
*
bn2blob_ptr
)[
"model_diffs"
]
=
KTCommon
::
CreateBlobWithSameValue
(
dim_vec
,
2
);
(
*
bn2blob_ptr
)[
"model_expected"
]
=
KTCommon
::
CreateBlobWithSameValue
(
dim_vec
,
2
);
(
*
bn2blob_ptr
)[
"mean_square_expected"
]
=
KTCommon
::
CreateBlobWithSameValue
(
dim_vec
,
3
);
return
[
bn2blob_ptr
](
const
std
::
string
&
bn
)
{
return
bn2blob_ptr
->
at
(
bn
);
};
}
template
<
DeviceType
device_type
,
typename
FloatingPointType
>
void
TestRMSPropMdUpdateKernel
()
{
using
KTCommon
=
KernelTestCommon
<
device_type
,
FloatingPointType
>
;
KernelCtx
ctx
;
KTCommon
::
BuildKernelCtx
(
&
ctx
);
const
float
learning_rate
=
{
2.0
f
};
const
float
decay_rate
=
{
0.5
f
};
const
float
epsilon
=
{
1.0
f
};
auto
BnInOp2BlobPtr
=
BuildBnInOp2BlobPtr
<
device_type
,
FloatingPointType
>
();
auto
rmsprop_md_update_kernel
=
BuildRMSPropMdUpdateKernel
<
device_type
,
FloatingPointType
>
(
learning_rate
,
decay_rate
,
epsilon
);
int32_t
piece_size
=
1
;
int32_t
num_of_pieces_in_batch
=
2
;
InitJobDesc
(
piece_size
,
num_of_pieces_in_batch
);
rmsprop_md_update_kernel
->
Forward
(
ctx
,
BnInOp2BlobPtr
);
KTCommon
::
SyncStream
(
&
ctx
);
KTCommon
::
CheckResult
(
BnInOp2BlobPtr
,
"mean_square"
,
"mean_square_expected"
);
KTCommon
::
CheckResult
(
BnInOp2BlobPtr
,
"model"
,
"model_expected"
);
}
}
// namespace
}
// namespace test
TEST
(
RMSPropMdUpdateKernel
,
model_update_cpu
)
{
test
::
TestRMSPropMdUpdateKernel
<
DeviceType
::
kCPU
,
float
>
();
test
::
TestRMSPropMdUpdateKernel
<
DeviceType
::
kCPU
,
double
>
();
}
TEST
(
RMSPropMdUpdateKernel
,
model_update_gpu
)
{
test
::
TestRMSPropMdUpdateKernel
<
DeviceType
::
kGPU
,
float
>
();
test
::
TestRMSPropMdUpdateKernel
<
DeviceType
::
kGPU
,
double
>
();
}
}
// namespace oneflow
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录