Skip to content
项目
群组
代码片段
帮助
当前项目
正在载入...
登录 / 注册
切换导航面板
P
pytensor
项目
项目
详情
活动
周期分析
仓库
仓库
文件
提交
分支
标签
贡献者
图表
比较
统计图
议题
0
议题
0
列表
看板
标记
里程碑
合并请求
0
合并请求
0
CI / CD
CI / CD
流水线
作业
日程
统计图
Wiki
Wiki
代码片段
代码片段
成员
成员
折叠边栏
关闭边栏
活动
图像
聊天
创建新问题
作业
提交
问题看板
Open sidebar
testgroup
pytensor
Commits
bbd9ff3f
提交
bbd9ff3f
authored
2月 02, 2016
作者:
Harm de Vries
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
flake8
上级
5c172018
隐藏空白字符变更
内嵌
并排
正在显示
2 个修改的文件
包含
38 行增加
和
43 行删除
+38
-43
dnn.py
theano/sandbox/cuda/dnn.py
+30
-37
test_dnn.py
theano/sandbox/cuda/tests/test_dnn.py
+8
-6
没有找到文件。
theano/sandbox/cuda/dnn.py
浏览文件 @
bbd9ff3f
...
@@ -1351,7 +1351,7 @@ class GpuDnnPoolDesc(GpuOp):
...
@@ -1351,7 +1351,7 @@ class GpuDnnPoolDesc(GpuOp):
class
GpuDnnPool
(
DnnBase
):
class
GpuDnnPool
(
DnnBase
):
"""
"""
Pooling.
Pooling.
Parameters
Parameters
...
@@ -1371,7 +1371,7 @@ class GpuDnnPool(DnnBase):
...
@@ -1371,7 +1371,7 @@ class GpuDnnPool(DnnBase):
"""
"""
__props__
=
(
"mode"
,)
__props__
=
(
"mode"
,)
def
__init__
(
self
,
mode
=
'max'
):
def
__init__
(
self
,
mode
=
'max'
):
super
(
GpuDnnPool
,
self
)
.
__init__
()
super
(
GpuDnnPool
,
self
)
.
__init__
()
if
mode
==
'average'
:
if
mode
==
'average'
:
...
@@ -1382,13 +1382,13 @@ class GpuDnnPool(DnnBase):
...
@@ -1382,13 +1382,13 @@ class GpuDnnPool(DnnBase):
def
make_node
(
self
,
img
,
ws
,
stride
,
pad
):
def
make_node
(
self
,
img
,
ws
,
stride
,
pad
):
img
=
as_cuda_ndarray_variable
(
img
)
img
=
as_cuda_ndarray_variable
(
img
)
assert
(
img
.
ndim
in
[
4
,
5
])
assert
(
img
.
ndim
in
[
4
,
5
])
ws
=
tensor
.
as_tensor_variable
(
ws
)
ws
=
tensor
.
as_tensor_variable
(
ws
)
stride
=
tensor
.
as_tensor_variable
(
stride
)
stride
=
tensor
.
as_tensor_variable
(
stride
)
pad
=
tensor
.
as_tensor_variable
(
pad
)
pad
=
tensor
.
as_tensor_variable
(
pad
)
assert
ws
.
type
.
ndim
==
stride
.
type
.
ndim
and
ws
.
type
.
ndim
==
pad
.
type
.
ndim
assert
ws
.
type
.
ndim
==
stride
.
type
.
ndim
and
ws
.
type
.
ndim
==
pad
.
type
.
ndim
assert
ws
.
type
.
ndim
==
1
assert
ws
.
type
.
ndim
==
1
return
Apply
(
self
,
[
img
,
ws
,
stride
,
pad
],
[
img
.
type
()])
return
Apply
(
self
,
[
img
,
ws
,
stride
,
pad
],
[
img
.
type
()])
def
infer_shape
(
self
,
node
,
shape
):
def
infer_shape
(
self
,
node
,
shape
):
...
@@ -1457,7 +1457,7 @@ if (pool%(name)s != NULL) { cudnnDestroyPoolingDescriptor(pool%(name)s); }
...
@@ -1457,7 +1457,7 @@ if (pool%(name)s != NULL) { cudnnDestroyPoolingDescriptor(pool%(name)s); }
raise
Exception
(
"cudnn v1 do not support average_exc_pad"
)
raise
Exception
(
"cudnn v1 do not support average_exc_pad"
)
else
:
else
:
raise
NotImplementedError
(
"Unsupported pooling model."
)
raise
NotImplementedError
(
"Unsupported pooling model."
)
return
"""
return
"""
fprintf(stderr, "test_forward
\\
n");
fprintf(stderr, "test_forward
\\
n");
cudnnStatus_t err;
cudnnStatus_t err;
...
@@ -1530,8 +1530,8 @@ if (err != CUDNN_STATUS_SUCCESS) {
...
@@ -1530,8 +1530,8 @@ if (err != CUDNN_STATUS_SUCCESS) {
"""
%
dict
(
out
=
out
,
fail
=
sub
[
'fail'
],
"""
%
dict
(
out
=
out
,
fail
=
sub
[
'fail'
],
name
=
name
,
input
=
inputs
[
0
],
name
=
name
,
input
=
inputs
[
0
],
ws
=
ws
,
pad
=
pad
,
str
=
stride
,
ws
=
ws
,
pad
=
pad
,
str
=
stride
,
nd
=
node
.
inputs
[
0
]
.
ndim
-
2
,
input_desc
=
"input"
+
name
,
nd
=
node
.
inputs
[
0
]
.
ndim
-
2
,
input_desc
=
"input"
+
name
,
output_desc
=
"output"
+
name
,
output_desc
=
"output"
+
name
,
mode_flag
=
mode_flag
)
mode_flag
=
mode_flag
)
def
grad
(
self
,
inp
,
grads
):
def
grad
(
self
,
inp
,
grads
):
...
@@ -1550,8 +1550,8 @@ if (err != CUDNN_STATUS_SUCCESS) {
...
@@ -1550,8 +1550,8 @@ if (err != CUDNN_STATUS_SUCCESS) {
# not connected to desc
# not connected to desc
return
[[
1
],
[
0
],
[
0
],
[
0
]]
return
[[
1
],
[
0
],
[
0
],
[
0
]]
#
def c_code_cache_version(self):
def
c_code_cache_version
(
self
):
#
return (8, version())
return
(
8
,
version
())
class
GpuDnnPoolGrad
(
DnnBase
):
class
GpuDnnPoolGrad
(
DnnBase
):
...
@@ -1579,7 +1579,7 @@ class GpuDnnPoolGrad(DnnBase):
...
@@ -1579,7 +1579,7 @@ class GpuDnnPoolGrad(DnnBase):
"""
"""
__props__
=
(
'mode'
,)
__props__
=
(
'mode'
,)
def
__init__
(
self
,
mode
=
'max'
):
def
__init__
(
self
,
mode
=
'max'
):
super
(
GpuDnnPoolGrad
,
self
)
.
__init__
()
super
(
GpuDnnPoolGrad
,
self
)
.
__init__
()
if
mode
==
'average'
:
if
mode
==
'average'
:
...
@@ -1594,13 +1594,13 @@ class GpuDnnPoolGrad(DnnBase):
...
@@ -1594,13 +1594,13 @@ class GpuDnnPoolGrad(DnnBase):
assert
(
inp_grad
.
ndim
in
[
4
,
5
])
assert
(
inp_grad
.
ndim
in
[
4
,
5
])
out
=
as_cuda_ndarray_variable
(
out
)
out
=
as_cuda_ndarray_variable
(
out
)
assert
(
out
.
ndim
in
[
4
,
5
])
assert
(
out
.
ndim
in
[
4
,
5
])
ws
=
tensor
.
as_tensor_variable
(
ws
)
ws
=
tensor
.
as_tensor_variable
(
ws
)
stride
=
tensor
.
as_tensor_variable
(
stride
)
stride
=
tensor
.
as_tensor_variable
(
stride
)
pad
=
tensor
.
as_tensor_variable
(
pad
)
pad
=
tensor
.
as_tensor_variable
(
pad
)
assert
ws
.
type
.
ndim
==
stride
.
type
.
ndim
and
ws
.
type
.
ndim
==
pad
.
type
.
ndim
assert
ws
.
type
.
ndim
==
stride
.
type
.
ndim
and
ws
.
type
.
ndim
==
pad
.
type
.
ndim
assert
ws
.
type
.
ndim
==
1
assert
ws
.
type
.
ndim
==
1
return
Apply
(
self
,
[
inp
,
out
,
inp_grad
,
ws
,
stride
,
pad
],
return
Apply
(
self
,
[
inp
,
out
,
inp_grad
,
ws
,
stride
,
pad
],
[
inp
.
type
()])
[
inp
.
type
()])
...
@@ -1661,19 +1661,14 @@ if (output%(name)s != NULL) { cudnnDestroyTensorDescriptor(output%(name)s); }
...
@@ -1661,19 +1661,14 @@ if (output%(name)s != NULL) { cudnnDestroyTensorDescriptor(output%(name)s); }
if (output_grad
%(name)
s != NULL) { cudnnDestroyTensorDescriptor(output_grad
%(name)
s); }
if (output_grad
%(name)
s != NULL) { cudnnDestroyTensorDescriptor(output_grad
%(name)
s); }
if (pool
%(name)
s != NULL) { cudnnDestroyPoolingDescriptor(pool
%(name)
s); }
if (pool
%(name)
s != NULL) { cudnnDestroyPoolingDescriptor(pool
%(name)
s); }
"""
%
dict
(
name
=
name
)
"""
%
dict
(
name
=
name
)
# def perform(self, node, inputs_storage, output_storage):
# output_storage[0][0] = inputs_storage[0].copy()
# return
def
c_code
(
self
,
node
,
name
,
inputs
,
outputs
,
sub
):
def
c_code
(
self
,
node
,
name
,
inputs
,
outputs
,
sub
):
# raise NotImplementedError()
# Here the name out and inp are based on the cudnn definition.
# Here the name out and inp are based on the cudnn definition.
# Not the definition of this class.
# Not the definition of this class.
# This make it complicated.
# This make it complicated.
out
,
inp
,
inp_grad
,
ws
,
stride
,
pad
=
inputs
out
,
inp
,
inp_grad
,
ws
,
stride
,
pad
=
inputs
out_grad
,
=
outputs
out_grad
,
=
outputs
if
self
.
mode
==
'max'
:
if
self
.
mode
==
'max'
:
mode_flag
=
'CUDNN_POOLING_MAX'
mode_flag
=
'CUDNN_POOLING_MAX'
elif
self
.
mode
==
"average_inc_pad"
:
elif
self
.
mode
==
"average_inc_pad"
:
...
@@ -1770,16 +1765,15 @@ if (err%(name)s != CUDNN_STATUS_SUCCESS) {
...
@@ -1770,16 +1765,15 @@ if (err%(name)s != CUDNN_STATUS_SUCCESS) {
"""
%
dict
(
output_grad
=
out_grad
,
"""
%
dict
(
output_grad
=
out_grad
,
fail
=
sub
[
'fail'
],
name
=
name
,
fail
=
sub
[
'fail'
],
name
=
name
,
input
=
inp
,
input_grad
=
inp_grad
,
output
=
out
,
input
=
inp
,
input_grad
=
inp_grad
,
output
=
out
,
input_desc
=
"input"
+
name
,
input_desc
=
"input"
+
name
,
input_grad_desc
=
"input_grad"
+
name
,
input_grad_desc
=
"input_grad"
+
name
,
output_desc
=
"output"
+
name
,
output_desc
=
"output"
+
name
,
output_grad_desc
=
"output_grad"
+
name
,
output_grad_desc
=
"output_grad"
+
name
,
mode_flag
=
mode_flag
,
nd
=
node
.
inputs
[
0
]
.
ndim
-
2
,
mode_flag
=
mode_flag
,
nd
=
node
.
inputs
[
0
]
.
ndim
-
2
,
ws
=
ws
,
pad
=
pad
,
str
=
stride
)
ws
=
ws
,
pad
=
pad
,
str
=
stride
)
def
c_code_cache_version
(
self
):
def
c_code_cache_version
(
self
):
return
return
(
8
,
version
())
#return (7, version())
def
infer_shape
(
self
,
node
,
shape
):
def
infer_shape
(
self
,
node
,
shape
):
return
[
shape
[
0
]]
return
[
shape
[
0
]]
...
@@ -1801,7 +1795,7 @@ def dnn_pool(img, ws, stride=(1, 1), mode='max', pad=(0, 0)):
...
@@ -1801,7 +1795,7 @@ def dnn_pool(img, ws, stride=(1, 1), mode='max', pad=(0, 0)):
stride
stride
Subsampling stride (default: (1, 1)).
Subsampling stride (default: (1, 1)).
mode : {'max', 'average_inc_pad', 'average_exc_pad}
mode : {'max', 'average_inc_pad', 'average_exc_pad}
pad :
pad :
(pad_h, pad_w) padding information.
(pad_h, pad_w) padding information.
pad_h is the number of zero-valued pixels added to each of the top and
pad_h is the number of zero-valued pixels added to each of the top and
bottom borders.
bottom borders.
...
@@ -2296,11 +2290,11 @@ if True:
...
@@ -2296,11 +2290,11 @@ if True:
return
return
inp
,
out
,
inp_grad
=
node
.
inputs
inp
,
out
,
inp_grad
=
node
.
inputs
ds
=
node
.
op
.
ds
ds
=
node
.
op
.
ds
return
[
GpuDnnPoolGrad
(
mode
=
'max'
)(
gpu_contiguous
(
inp
),
return
[
GpuDnnPoolGrad
(
mode
=
'max'
)(
gpu_contiguous
(
inp
),
gpu_contiguous
(
out
),
gpu_contiguous
(
out
),
gpu_contiguous
(
inp_grad
),
gpu_contiguous
(
inp_grad
),
ds
,
ds
,
(
0
,
0
))]
ds
,
ds
,
(
0
,
0
))]
@register_opt
(
'cudnn'
)
@register_opt
(
'cudnn'
)
@local_optimizer
([
MaxPoolGrad
])
@local_optimizer
([
MaxPoolGrad
])
...
@@ -2322,9 +2316,9 @@ if True:
...
@@ -2322,9 +2316,9 @@ if True:
HostFromGpu
))):
HostFromGpu
))):
ret
=
GpuDnnPoolGrad
(
mode
=
mode
)(
gpu_contiguous
(
inp
),
ret
=
GpuDnnPoolGrad
(
mode
=
mode
)(
gpu_contiguous
(
inp
),
gpu_contiguous
(
out
),
gpu_contiguous
(
out
),
gpu_contiguous
(
inp_grad
),
gpu_contiguous
(
inp_grad
),
ds
,
st
,
pad
)
ds
,
st
,
pad
)
return
[
host_from_gpu
(
ret
)]
return
[
host_from_gpu
(
ret
)]
@register_opt
(
'cudnn'
)
@register_opt
(
'cudnn'
)
...
@@ -2346,12 +2340,11 @@ if True:
...
@@ -2346,12 +2340,11 @@ if True:
HostFromGpu
))):
HostFromGpu
))):
contiguous_inp_grad
=
gpu_contiguous
(
inp_grad
)
contiguous_inp_grad
=
gpu_contiguous
(
inp_grad
)
ret
=
GpuDnnPoolGrad
(
mode
=
mode
)(
gpu_contiguous
(
inp
),
ret
=
GpuDnnPoolGrad
(
mode
=
mode
)(
gpu_contiguous
(
inp
),
contiguous_inp_grad
,
contiguous_inp_grad
,
contiguous_inp_grad
,
contiguous_inp_grad
,
ds
,
st
,
pad
)
ds
,
st
,
pad
)
return
[
host_from_gpu
(
ret
)]
return
[
host_from_gpu
(
ret
)]
@register_opt
(
'cudnn'
)
@register_opt
(
'cudnn'
)
@local_optimizer
([
GpuSoftmax
])
@local_optimizer
([
GpuSoftmax
])
def
local_softmax_dnn
(
node
):
def
local_softmax_dnn
(
node
):
...
...
theano/sandbox/cuda/tests/test_dnn.py
浏览文件 @
bbd9ff3f
...
@@ -240,7 +240,7 @@ def test_pooling():
...
@@ -240,7 +240,7 @@ def test_pooling():
modes
=
(
'max'
,
'average_inc_pad'
)
modes
=
(
'max'
,
'average_inc_pad'
)
else
:
else
:
modes
=
(
'max'
,
'average_inc_pad'
,
'average_exc_pad'
)
modes
=
(
'max'
,
'average_inc_pad'
,
'average_exc_pad'
)
x
=
T
.
ftensor4
()
x
=
T
.
ftensor4
()
for
mode
,
pad
in
product
(
modes
,
for
mode
,
pad
in
product
(
modes
,
((
0
,
0
),
(
1
,
0
),
(
1
,
0
),
(
2
,
3
),
(
3
,
2
))):
((
0
,
0
),
(
1
,
0
),
(
1
,
0
),
(
2
,
3
),
(
3
,
2
))):
...
@@ -300,7 +300,8 @@ def test_pooling():
...
@@ -300,7 +300,8 @@ def test_pooling():
# Not implemented
# Not implemented
continue
continue
pad_
=
theano
.
shared
(
numpy
.
array
(
pad
))
pad_
=
theano
.
shared
(
numpy
.
array
(
pad
))
## This test the CPU grad + opt + GPU implemtentation
# This test the CPU grad + opt + GPU implemtentation
def
fn
(
x
):
def
fn
(
x
):
return
pool_2d
(
x
,
(
2
,
2
),
ignore_border
=
True
,
return
pool_2d
(
x
,
(
2
,
2
),
ignore_border
=
True
,
padding
=
pad
,
mode
=
mode
)
padding
=
pad
,
mode
=
mode
)
...
@@ -312,7 +313,7 @@ def test_pooling():
...
@@ -312,7 +313,7 @@ def test_pooling():
mode
=
mode_with_gpu
)
mode
=
mode_with_gpu
)
assert
any
([
isinstance
(
node
.
op
,
cuda
.
dnn
.
GpuDnnPoolGrad
)
assert
any
([
isinstance
(
node
.
op
,
cuda
.
dnn
.
GpuDnnPoolGrad
)
for
node
in
fg
.
maker
.
fgraph
.
toposort
()])
for
node
in
fg
.
maker
.
fgraph
.
toposort
()])
# Test the GPU grad + GPU implementation
# Test the GPU grad + GPU implementation
def
fn
(
x
):
def
fn
(
x
):
dnn_op
=
cuda
.
dnn
.
dnn_pool
(
dnn_op
=
cuda
.
dnn
.
dnn_pool
(
...
@@ -321,7 +322,7 @@ def test_pooling():
...
@@ -321,7 +322,7 @@ def test_pooling():
pad
=
pad_
,
pad
=
pad_
,
mode
=
mode
)
mode
=
mode
)
return
dnn_op
return
dnn_op
theano
.
tests
.
unittest_tools
.
verify_grad
(
theano
.
tests
.
unittest_tools
.
verify_grad
(
fn
,
[
data
],
fn
,
[
data
],
cast_to_output_type
=
False
,
cast_to_output_type
=
False
,
...
@@ -337,7 +338,7 @@ def test_pooling():
...
@@ -337,7 +338,7 @@ def test_pooling():
out
=
pool_2d
(
x
,
(
2
,
2
),
st
=
(
1
,
1
),
out
=
pool_2d
(
x
,
(
2
,
2
),
st
=
(
1
,
1
),
padding
=
pad
,
padding
=
pad
,
ignore_border
=
True
,
mode
=
mode
)
ignore_border
=
True
,
mode
=
mode
)
fc
=
theano
.
function
([
x
],
theano
.
grad
(
out
.
sum
(),
x
),
fc
=
theano
.
function
([
x
],
theano
.
grad
(
out
.
sum
(),
x
),
mode
=
mode_without_gpu
)
mode
=
mode_without_gpu
)
if
mode
==
'max'
:
if
mode
==
'max'
:
...
@@ -1005,7 +1006,8 @@ class TestDnnInferShapes(utt.InferShapeTester):
...
@@ -1005,7 +1006,8 @@ class TestDnnInferShapes(utt.InferShapeTester):
):
):
self
.
_compile_and_check
(
self
.
_compile_and_check
(
[
img
],
[
img
],
[
dnn
.
GpuDnnPool
(
mode
=
params
[
2
])(
img
,
params
[
0
],
params
[
1
],
(
0
,
0
))],
[
dnn
.
GpuDnnPool
(
mode
=
params
[
2
])
(
img
,
params
[
0
],
params
[
1
],
(
0
,
0
))],
[
img_val
],
[
img_val
],
dnn
.
GpuDnnPool
dnn
.
GpuDnnPool
)
)
...
...
编写
预览
Markdown
格式
0%
重试
或
添加新文件
添加附件
取消
您添加了
0
人
到此讨论。请谨慎行事。
请先完成此评论的编辑!
取消
请
注册
或者
登录
后发表评论