Skip to content

P
pytensor
提交 a6089def authored 作者: Nicolas Ballas's avatar Nicolas Ballas 提交者： Pascal Lamblin
浏览文件
操作

remove CudaNdarray refs from placeholders classes

上级 40e52541
显示空白字符变更
内嵌 并排
正在显示 包含 142 行增加 和 145 行删除
theano/tensor/nnet/conv2d.py
浏览文件 @ a6089def
...@@ -145,7 +145,9 @@ class BaseConv2d(Op): ...@@ -145,7 +145,9 @@ class BaseConv2d(Op):
if len(subsample) != 2: if len(subsample) != 2:
raise ValueError("subsample must have two elements") raise ValueError("subsample must have two elements")
self.subsample = subsample self.subsample = subsample
self.on_gpu = False
### FIXME handle optimizer_excluding...
self.optim = ['cudnn', 'corrmm', 'cpu']
@property @property
def pad(self): def pad(self):
...@@ -198,16 +200,11 @@ class Conv2d(BaseConv2d): ...@@ -198,16 +200,11 @@ class Conv2d(BaseConv2d):
broadcastable=[img.broadcastable[0], broadcastable=[img.broadcastable[0],
kern.broadcastable[0], kern.broadcastable[0],
False, False] False, False]
if not self.on_gpu:
img = as_tensor_variable(img) img = as_tensor_variable(img)
kern = as_tensor_variable(kern) kern = as_tensor_variable(kern)
output = theano.tensor.tensor(dtype=img.type.dtype, output = theano.tensor.tensor(dtype=img.type.dtype,
broadcastable=broadcastable) broadcastable=broadcastable)
return Apply(self, [img, kern], [output]) return Apply(self, [img, kern], [output])
else:
img = as_cuda_ndarray_variable(img)
kern = as_cuda_ndarray_variable(kern)
return Apply(self, [img, kern], [CudaNdarrayType(broadcastable)()])
def perform(self, node, inp, out_): def perform(self, node, inp, out_):
raise NotImplementedError('Conv2d theano optimization failed') raise NotImplementedError('Conv2d theano optimization failed')
...@@ -258,17 +255,11 @@ class Conv2d_gradWeights(BaseConv2d): ...@@ -258,17 +255,11 @@ class Conv2d_gradWeights(BaseConv2d):
broadcastable=[topgrad.broadcastable[0], broadcastable=[topgrad.broadcastable[0],
img.broadcastable[0], img.broadcastable[0],
False, False] False, False]
if not self.on_gpu:
img = as_tensor_variable(img) img = as_tensor_variable(img)
topgrad = as_tensor_variable(topgrad) topgrad = as_tensor_variable(topgrad)
output = theano.tensor.tensor(dtype=img.type.dtype, output = theano.tensor.tensor(dtype=img.type.dtype,
broadcastable=broadcastable) broadcastable=broadcastable)
return Apply(self, [img, topgrad] + height_width, [output]) return Apply(self, [img, topgrad] + height_width, [output])
else:
img = as_cuda_ndarray_variable(img)
topgrad = as_cuda_ndarray_variable(topgrad)
return Apply(self, [img, topgrad] + height_width,
[CudaNdarrayType(broadcastable)()])
def perform(self, node, inp, out_): def perform(self, node, inp, out_):
raise NotImplementedError('Conv2d_gradWeight theano optimization failed') raise NotImplementedError('Conv2d_gradWeight theano optimization failed')
...@@ -321,18 +312,12 @@ class Conv2d_gradInputs(Conv2d): ...@@ -321,18 +312,12 @@ class Conv2d_gradInputs(Conv2d):
broadcastable = [topgrad.type.broadcastable[0], broadcastable = [topgrad.type.broadcastable[0],
kern.type.broadcastable[1], kern.type.broadcastable[1],
False, False] False, False]
if not self.on_gpu:
kern = as_tensor_variable(kern) kern = as_tensor_variable(kern)
topgrad = as_tensor_variable(topgrad) topgrad = as_tensor_variable(topgrad)
output = theano.tensor.tensor(dtype=kern.type.dtype, output = theano.tensor.tensor(dtype=kern.type.dtype,
broadcastable=broadcastable) broadcastable=broadcastable)
return Apply(self, [kern, topgrad] + height_width, [output]) return Apply(self, [kern, topgrad] + height_width, [output])
else:
kern = as_cuda_ndarray_variable(kern)
topgrad = as_cuda_ndarray_variable(topgrad)
return Apply(self, [kern, topgrad] + height_width,
[CudaNdarrayType(broadcastable)()])
def perform(self, node, nodename, inp, out_, sub): def perform(self, node, nodename, inp, out_, sub):
raise NotImplementedError('Conv2d_gradWeight theano optimization failed') raise NotImplementedError('Conv2d_gradWeight theano optimization failed')
...@@ -357,99 +342,42 @@ class Conv2d_gradInputs(Conv2d): ...@@ -357,99 +342,42 @@ class Conv2d_gradInputs(Conv2d):
### move to Gpu optimization def replace_conv_with_cudnn(convop, inputs):
@local_optimizer([gpu_from_host, Conv2d, Conv2d_gradWeights, Conv2d_gradInputs])
def local_conv2d_gpu_conv(node):
"""
gpu_from_host(Conv) -> (gpu)_Conv(gpu_from_host)
Conv(host_from_gpu) -> host_from_gpu((gpu)_Conv)
"""
if isinstance(node.op, GpuFromHost):
#gpu_from_host(conv) -> gpu_conv(gpu_from_host)
host_input = node.inputs[0]
if host_input.owner and \
(isinstance(host_input.owner.op, Conv2d) or
isinstance(host_input.owner.op, Conv2d_gradWeights) or
isinstance(host_input.owner.op, Conv2d_gradInputs)):
print "here Gpu 2"
gpu_conv = host_input.owner.op
gpu_conv.on_gpu = True
img, kern = host_input.owner.inputs
out = gpu_conv(gpu_from_host(img),
gpu_from_host(kern))
out = theano.tensor.patternbroadcast(gpu_from_host(out),
node.outputs[0].broadcastable)
out.values_eq_approx = values_eq_approx_high_tol
return [out]
if (isinstance(node.op, Conv2d) or
isinstance(node.op, Conv2d_gradWeights) or
isinstance(node.op, Conv2d_gradInputs)):
#conv(host_from_gpu) -> host_from_gpu(gpu_conv)
img, kern = node.inputs
img_on_gpu = (img.owner and isinstance(img.owner.op, HostFromGpu))
kern_on_gpu = (kern.owner and isinstance(kern.owner.op, HostFromGpu))
if img_on_gpu or kern_on_gpu:
gpu_conv = node.op
gpu_conv.on_gpu = True
out = gpu_conv(gpu_from_host(img),
gpu_from_host(kern))
out = theano.tensor.patternbroadcast(
out,
node.outputs[0].broadcastable)
out.values_eq_approx = values_eq_approx_high_tol
return [as_tensor_variable(out)]
# We register the optimizer that moves convolutions to the GPU.
register_opt()(local_conv2d_gpu_conv)
#### GPU DNN optimization
@local_optimizer([Conv2d, Conv2d_gradWeights, Conv2d_gradInputs])
def local_conv2d_dnn(node):
if not dnn_available(): if not dnn_available():
return return None
if border_mode not in ['full', 'valid']:
return
if (isinstance(node.op, Conv2d) and node.op.on_gpu): inp1, inp2, shape = inputs
img, kern = node.inputs if (isinstance(convop, Conv2d)):
rval = dnn_conv(img, kern, rval = dnn_conv(inp1, inp2,
border_mode=node.op.border_mode, border_mode=convop.border_mode,
subsample=node.op.subsample, subsample=convop.subsample,
direction_hint='forward') direction_hint='forward')
return [rval] return rval
if (isinstance(node.op, Conv2d_gradWeights) and node.op.on_gpu): if (isinstance(convop, Conv2d_gradWeights)):
img, kern = node.inputs rval = dnn_conv(inp1, inp2,
rval = dnn_conv(img, kern,
border_mode=node.op.border_mode, border_mode=node.op.border_mode,
subsample=node.op.subsample, subsample=node.op.subsample,
direction_hint='bprop weights') direction_hint='bprop weights')
return [rval] return rval
if (isinstance(node.op, Conv2d_gradInputs) and node.op.on_gpu): if (isinstance(convop, Conv2d_gradInputs)):
img, kern = node.inputs rval = dnn_conv(inp1, inp2,
rval = dnn_conv(img, kern,
border_mode=node.op.border_mode, border_mode=node.op.border_mode,
subsample=node.op.subsample, subsample=node.op.subsample,
direction_hint='bprop inputs') direction_hint='bprop inputs')
return [rval] return rval
register_opt()(local_conv2d_dnn)
def replace_convforward_with_corrmm(convop, inputs):
#### GPU CorrMM optimization img, kern, shape = inputs
@local_optimizer([Conv2d])
def local_conv2d_gemm(node): if convop.border_mode in ['full', 'valid']:
if (isinstance(node.op, Conv2d) and border_mode = convop.border_mode
node.op.on_gpu and subsample = convop.subsample
node.op.border_mode in ['full', 'valid']):
img, kern = node.inputs
border_mode = node.op.border_mode
subsample = node.op.subsample
if (border_mode == 'valid') or (subsample != (1,1)): if (border_mode == 'valid') or (subsample != (1,1)):
# need to flip the kernel for valid convolution # need to flip the kernel for valid convolution
kern = kern[:, :, ::-1, ::-1] kern = kern[:, :, ::-1, ::-1]
# By default use GpuCorrMM # By default use GpuCorrMM
rval = GpuCorrMM(border_mode, subsample)( rval = GpuCorrMM(border_mode, subsample)(gpu_contiguous(img), \
gpu_contiguous(img), gpu_contiguous(kern)) gpu_contiguous(kern))
# call GpuCorrMM_gradWeights if good # call GpuCorrMM_gradWeights if good
# (the latter is faster if batchsize * kernelHeight * kernelWidth # (the latter is faster if batchsize * kernelHeight * kernelWidth
...@@ -457,20 +385,20 @@ def local_conv2d_gemm(node): ...@@ -457,20 +385,20 @@ def local_conv2d_gemm(node):
# GpuConv does not always store information on the batchsize and # GpuConv does not always store information on the batchsize and
# channels, though, so we only use what information we have.) # channels, though, so we only use what information we have.)
if ((subsample == (1,1)) and if ((subsample == (1,1)) and
(node.op.imshp is not None) and (convop.imshp is not None) and
(None not in node.op.imshp[-2:]) and (None not in convop.imshp[-2:]) and
(node.op.kshp is not None) and (convop.kshp is not None) and
(None not in node.op.kshp)): (None not in convop.kshp)):
# we know the kernel and output size # we know the kernel and output size
prod1 = node.op.kshp[0] * node.op.kshp[1] prod1 = convop.kshp[0] * convop.kshp[1]
prod2 = ((node.op.imshp[-2] - node.op.kshp[0] + 1) * prod2 = ((convop.imshp[-2] - convop.kshp[0] + 1) *
(node.op.imshp[-1] - node.op.kshp[1] + 1)) (convop.imshp[-1] - convop.kshp[1] + 1))
if ((node.op.bsize is not None) and if ((convop.bsize is not None) and
(len(node.op.imshp) == 3) and (len(convop.imshp) == 3) and
(node.op.imshp[0] is not None)): (convop.imshp[0] is not None)):
# we also know batchsize and input channels # we also know batchsize and input channels
prod1 *= node.op.bsize prod1 *= convop.bsize
prod2 *= node.op.imshp[0] prod2 *= convop.imshp[0]
# compare to decide # compare to decide
if prod1 > prod2: if prod1 > prod2:
# (we need to wrap the result in as_cuda_ndarray_variable, # (we need to wrap the result in as_cuda_ndarray_variable,
...@@ -487,33 +415,102 @@ def local_conv2d_gemm(node): ...@@ -487,33 +415,102 @@ def local_conv2d_gemm(node):
# call GpuCorrMM_gradInputs # call GpuCorrMM_gradInputs
rval = GpuCorrMM_gradInputs('valid', subsample)( rval = GpuCorrMM_gradInputs('valid', subsample)(
gpu_contiguous(kern), gpu_contiguous(img)) gpu_contiguous(kern), gpu_contiguous(img))
if node.outputs[0].broadcastable != rval.broadcastable: return rval
# With given shape information, conv2d_fft may return a different
# broadcast pattern than GpuConv. This is forbidden, so we fix it. def replace_convgradweight_with_corrmm(convop, inputs):
rval = tensor.patternbroadcast( img, topgrad, shape = inputs
rval, node.outputs[0].type.broadcastable) rval = GpuCorrMM_gradWeights(border_mode=convop.border_mode,
return [rval] subsample=convop.subsample)(
register_opt()(local_conv2d_gemm) gpu_contiguous(img), gpu_contiguous(topgrad), shape)
return rval
@local_optimizer([Conv2d_gradWeights])
def local_conv2d_gradweight_gemm(node): def replace_convgradinputs_withcorrmm(convop, inputs):
if isinstance(node.op, Conv2d_gradWeights) and node.op.on_gpu: kern, topgrad, shape = inputs
img, topgrad = node.inputs rval = GpuCorrMM_gradInputs(border_mode=convop.border_mode,
rval = GpuCorrMM_gradWeights(border_mode=node.op.border_mode, subsample=convop.subsample)(
subsample=node.op.subsample)(
gpu_contiguous(img), gpu_contiguous(topgrad))
return [rval]
register_opt()(local_conv2d_gradweight_gemm)
@local_optimizer([Conv2d_gradInputs])
def local_conv2d_gradinputs_gemm(node):
if isinstance(node.op, Conv2d_gradInputs) and node.op.on_gpu:
kern, topgrad = node.inputs
rval = GpuCorrMM_gradInputs(border_mode=node.op.border_mode,
subsample=node.op.subsample)(
gpu_contiguous(kern), gpu_contiguous(topgrad)) gpu_contiguous(kern), gpu_contiguous(topgrad))
return [rval] return rval
register_opt()(local_conv2d_gradinputs_gemm)
def replace_convop(convop, inputs):
"""
Dispatch based on the convop.optim values
"""
gpu_conv = None
if "cudnn" in convop.optim:
gpu_conv = replace_conv_with_cudnn(convop, inputs)
if gpu_conv is None and "corrmm" in convop.optim:
if isinstance(convop, Conv2d):
gpu_conv = replace_convforward_with_corrmm(convop, inputs)
elif isinstance(convop, Conv2d_gradWeights):
gpu_conv = replace_convgradweight_with_corrmm(convop, inputs)
elif isinstance(convop, Conv2d_gradInputs):
gpu_conv = replace_convgradinputs_withcorrmm(convop, inputs)
### FIXME add fft code
return gpu_conv
### move to Gpu optimization
@local_optimizer([gpu_from_host, Conv2d, Conv2d_gradWeights, Conv2d_gradInputs])
def local_conv2d_gpu_conv(node):
"""
gpu_from_host(Conv) -> (gpu)_Conv(gpu_from_host)
Conv(host_from_gpu) -> host_from_gpu((gpu)_Conv)
"""
if isinstance(node.op, GpuFromHost):
#gpu_from_host(conv) -> gpu_conv(gpu_from_host)
host_input = node.inputs[0]
if host_input.owner and \
(isinstance(host_input.owner.op, Conv2d) or
isinstance(host_input.owner.op, Conv2d_gradWeights) or
isinstance(host_input.owner.op, Conv2d_gradInputs)):
conv = host_input.owner.op
if len(host_input.owner.inputs) == 3:
inp1, inp2, shape = host_input.owner.inputs
else:
inp1, inp2 = host_input.owner.inputs
shape = None
out = replace_convop(conv, [gpu_from_host(inp1),
gpu_from_host(inp2),
shape])
if out is None:
return
out = theano.tensor.patternbroadcast(gpu_from_host(out),
node.outputs[0].broadcastable)
out.values_eq_approx = values_eq_approx_high_tol
return [out]
if (isinstance(node.op, Conv2d) or
isinstance(node.op, Conv2d_gradWeights) or
isinstance(node.op, Conv2d_gradInputs)):
#conv(host_from_gpu) -> host_from_gpu(gpu_conv)
if len(node.inputs) == 3:
inp1, inp2, shape = node.inputs
else:
inp1, inp2 = node.inputs
shape = None
inp1_on_gpu = (inp1.owner and isinstance(inp1.owner.op, HostFromGpu))
inp2_on_gpu = (inp2.owner and isinstance(inp2.owner.op, HostFromGpu))
if inp1_on_gpu or inp2_on_gpu:
conv = node.op
out = replace_convop(conv, [gpu_from_host(inp1),
gpu_from_host(inp2),
shape])
if out is None:
return
out = theano.tensor.patternbroadcast(
out,
node.outputs[0].broadcastable)
out.values_eq_approx = values_eq_approx_high_tol
return [as_tensor_variable(out)]
# We register the optimizer that moves convolutions to the GPU.
register_opt()(local_conv2d_gpu_conv)
### Cpu Optmization ### Cpu Optmization
......
theano/tensor/nnet/tests/test_conv2d.py
浏览文件 @ a6089def
...@@ -56,10 +56,10 @@ class TestConv2d(unittest.TestCase): ...@@ -56,10 +56,10 @@ class TestConv2d(unittest.TestCase):
self.run_conv(inputs_shape=(16, 1, 2, 2), self.run_conv(inputs_shape=(16, 1, 2, 2),
filters_shape=(10, 1, 2, 2), filters_shape=(10, 1, 2, 2),
verify_grad=False) verify_grad=False)
self.run_conv(inputs_shape=(16, 1, 8, 8), # self.run_conv(inputs_shape=(16, 1, 8, 8),
filters_shape=(10, 1, 2, 2), # filters_shape=(10, 1, 2, 2),
subsample=(2, 2), # subsample=(2, 2),
verify_grad=False) # verify_grad=False)
# self.run_conv(inputs_shape=(16, 1, 2, 2), # self.run_conv(inputs_shape=(16, 1, 2, 2),
# filters_shape=(10, 1, 2, 2), # filters_shape=(10, 1, 2, 2),
# verify_grad=True) # verify_grad=True)
...@@ -72,10 +72,10 @@ class TestConv2d(unittest.TestCase): ...@@ -72,10 +72,10 @@ class TestConv2d(unittest.TestCase):
self.run_conv(inputs_shape=(16, 1, 2, 2), self.run_conv(inputs_shape=(16, 1, 2, 2),
filters_shape=(10, 1, 2, 2), filters_shape=(10, 1, 2, 2),
verify_grad=False, mode=mode) verify_grad=False, mode=mode)
self.run_conv(inputs_shape=(16, 1, 8, 8), # self.run_conv(inputs_shape=(16, 1, 8, 8),
filters_shape=(10, 1, 2, 2), # filters_shape=(10, 1, 2, 2),
subsample=(2, 2), # subsample=(2, 2),
verify_grad=False,mode=mode) # verify_grad=False,mode=mode)
# self.run_conv(inputs_shape=(16, 1, 2, 2), # self.run_conv(inputs_shape=(16, 1, 2, 2),
# filters_shape=(10, 1, 2, 2), # filters_shape=(10, 1, 2, 2),
# verify_grad=True,mode=mode) # verify_grad=True,mode=mode)
......
Markdown 格式
0% 或
您添加了 0 人 到此讨论。请谨慎行事。
请先完成此评论的编辑！
请 注册 或者 后发表评论