Theano: OpenCL and Theano/Keras Support for GPU
I think similar issue as : https://github.com/Theano/Theano/issues/5504
GPU: Radeon Pro 560, OS X
Theano is running on GPU via OpenCL, not CUDA. I ran the script here "http://deeplearning.net/software/theano/tutorial/using_gpu.html"
However, I am getting the following GPUArrayException error when I use model.fit_generator
GpuArrayException Traceback (most recent call last)
76 epochs=epochs,
77 validation_data=validation_generator,
---> 78 validation_steps=nb_validation_samples // batch_size)
79
80
/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/legacy/interfaces.pyc in wrapper(args, *kwargs)
85 warnings.warn('Update your ' + object_name +
86 ' call to the Keras 2 API: ' + signature, stacklevel=2)
---> 87 return func(args, *kwargs)
88 wrapper._original_function = func
89 return wrapper
/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/models.pyc in fit_generator(self, generator, steps_per_epoch, epochs, verbose, callbacks, validation_data, validation_steps, class_weight, max_queue_size, workers, use_multiprocessing, initial_epoch)
1115 workers=workers,
1116 use_multiprocessing=use_multiprocessing,
-> 1117 initial_epoch=initial_epoch)
1118
1119 @interfaces.legacy_generator_methods_support
/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/legacy/interfaces.pyc in wrapper(args, *kwargs)
85 warnings.warn('Update your ' + object_name +
86 ' call to the Keras 2 API: ' + signature, stacklevel=2)
---> 87 return func(args, *kwargs)
88 wrapper._original_function = func
89 return wrapper
/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/engine/training.pyc in fit_generator(self, generator, steps_per_epoch, epochs, verbose, callbacks, validation_data, validation_steps, class_weight, max_queue_size, workers, use_multiprocessing, initial_epoch)
1723
1724 do_validation = bool(validation_data)
-> 1725 self._make_train_function()
1726 if do_validation:
1727 self._make_test_function()
/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/engine/training.pyc in _make_train_function(self)
942 updates=updates,
943 name='train_function',
--> 944 **self._function_kwargs)
945
946 def _make_test_function(self):
/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/backend/theano_backend.pyc in function(inputs, outputs, updates, *kwargs)
1204 msg = 'Invalid argument "%s" passed to K.function with Theano backend' % key
1205 raise ValueError(msg)
-> 1206 return Function(inputs, outputs, updates=updates, *kwargs)
1207
1208
/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/backend/theano_backend.pyc in __init__(self, inputs, outputs, updates, name, *kwargs)
1190 on_unused_input='ignore',
1191 name=name,
-> 1192 *kwargs)
1193 self.name = name
1194
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/compile/function.pyc in function(inputs, outputs, mode, updates, givens, no_default_updates, accept_inplace, name, rebuild_strict, allow_input_downcast, profile, on_unused_input)
324 on_unused_input=on_unused_input,
325 profile=profile,
--> 326 output_keys=output_keys)
327 # We need to add the flag check_aliased inputs if we have any mutable or
328 # borrowed used defined inputs
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/compile/pfunc.pyc in pfunc(params, outputs, mode, updates, givens, no_default_updates, accept_inplace, name, rebuild_strict, allow_input_downcast, profile, on_unused_input, output_keys)
484 accept_inplace=accept_inplace, name=name,
485 profile=profile, on_unused_input=on_unused_input,
--> 486 output_keys=output_keys)
487
488
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/compile/function_module.pyc in orig_function(inputs, outputs, mode, accept_inplace, name, profile, on_unused_input, output_keys)
1793 on_unused_input=on_unused_input,
1794 output_keys=output_keys).create(
-> 1795 defaults)
1796
1797 t2 = time.time()
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/compile/function_module.pyc in create(self, input_storage, trustme, storage_map)
1659 theano.config.traceback.limit = theano.config.traceback.compile_limit
1660 _fn, _i, _o = self.linker.make_thunk(
-> 1661 input_storage=input_storage_lists, storage_map=storage_map)
1662 finally:
1663 theano.config.traceback.limit = limit_orig
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/link.pyc in make_thunk(self, input_storage, output_storage, storage_map)
697 return self.make_all(input_storage=input_storage,
698 output_storage=output_storage,
--> 699 storage_map=storage_map)[:3]
700
701 def make_all(self, input_storage, output_storage):
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/vm.pyc in make_all(self, profiler, input_storage, output_storage, storage_map)
1045 compute_map,
1046 no_recycling,
-> 1047 impl=impl))
1048 linker_make_thunk_time[node] = time.time() - thunk_start
1049 if not hasattr(thunks[-1], 'lazy'):
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/op.pyc in make_thunk(self, node, storage_map, compute_map, no_recycling, impl)
933 try:
934 return self.make_c_thunk(node, storage_map, compute_map,
--> 935 no_recycling)
936 except (NotImplementedError, utils.MethodNotDefined):
937 # We requested the c code, so don't catch the error.
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/op.pyc in make_c_thunk(self, node, storage_map, compute_map, no_recycling)
837 _logger.debug('Trying CLinker.make_thunk')
838 outputs = cl.make_thunk(input_storage=node_input_storage,
--> 839 output_storage=node_output_storage)
840 fill_storage, node_input_filters, node_output_filters = outputs
841
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in make_thunk(self, input_storage, output_storage, storage_map, keep_lock)
1188 cthunk, in_storage, out_storage, error_storage = self.__compile__(
1189 input_storage, output_storage, storage_map,
-> 1190 keep_lock=keep_lock)
1191
1192 res = _CThunk(cthunk, init_tasks, tasks, error_storage)
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in __compile__(self, input_storage, output_storage, storage_map, keep_lock)
1129 output_storage,
1130 storage_map,
-> 1131 keep_lock=keep_lock)
1132 return (thunk,
1133 [link.Container(input, storage) for input, storage in
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in cthunk_factory(self, error_storage, in_storage, out_storage, storage_map, keep_lock)
1584 node.op.prepare_node(node, storage_map, None, 'c')
1585 module = get_module_cache().module_from_key(
-> 1586 key=key, lnk=self, keep_lock=keep_lock)
1587
1588 vars = self.inputs + self.outputs + self.orphans
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cmodule.pyc in module_from_key(self, key, lnk, keep_lock)
1120 return module
1121
-> 1122 src_code = lnk.get_src_code()
1123 # Is the source code already in the cache?
1124 module_hash = get_module_hash(src_code, key)
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in get_src_code(self)
1460
1461 def get_src_code(self):
-> 1462 mod = self.get_dynamic_module()
1463 return mod.code()
1464
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in get_dynamic_module(self)
1504 """
1505 if not hasattr(self, '_mod'):
-> 1506 self.code_gen()
1507
1508 mod = cmodule.DynamicModule()
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in code_gen(self)
780 try:
781 c_support_code_apply.append(op.c_support_code_apply(node,
--> 782 name))
783 except utils.MethodNotDefined:
784 pass
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gpuarray/basic_ops.pyc in f(args)
450 def forward_string_meth(name):
451 def f(args):
--> 452 res = getattr(GpuKernelBase, name)(args)
453 try:
454 res = res + '\n' + getattr(COp, name)(args)
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gpuarray/basic_ops.pyc in c_support_code_apply(self, node, name)
375 kernels = self.gpu_kernels(node, name)
376 ctx = self.get_params(node)
--> 377 bins = '\n'.join(self._generate_kernel_bin(k, ctx) for k in kernels)
378 codes = '\n'.join(self._generate_kernel_code(k) for k in kernels)
379 return '\n'.join([bins, codes])
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gpuarray/basic_ops.pyc in
375 kernels = self.gpu_kernels(node, name)
376 ctx = self.get_params(node)
--> 377 bins = '\n'.join(self._generate_kernel_bin(k, ctx) for k in kernels)
378 codes = '\n'.join(self._generate_kernel_code(k) for k in kernels)
379 return '\n'.join([bins, codes])
/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gpuarray/basic_ops.pyc in _generate_kernel_bin(self, k, ctx)
317 def _generate_kernel_bin(self, k, ctx):
318 gk = gpuarray.GpuKernel(k.code, k.name, k.params, context=ctx,
--> 319 **k._get_py_flags())
320 bin = gk._binary
321 bcode = ','.join(hex(c) for c in iterbytes(bin))
pygpu/gpuarray.pyx in pygpu.gpuarray.GpuKernel.__cinit__()
pygpu/gpuarray.pyx in pygpu.gpuarray.kernel_init()
GpuArrayException: ('The following error happened while compiling the node', GpuCorrMM{valid, (1, 1), (1, 1)}(GpuContiguous.0, GpuContiguous.0), '\n', u'Program build failure ::\n~~\n~~~~~\n~~ ^ ~\n~ ^ ~\n0001\t#define local_barrier() barrier(CLK_LOCAL_MEM_FENCE)\n0002\t#define WITHIN_KERNEL /
I dug around a bit and found out that in feb 2017 someone said that theano stopped openCL support. Is this still the case?
aditsanghvi94
All 3 comments
We didn't stopped OpenCL support, we never finished it. We don't have time
to finish it, so unless you have a significant time to help us finish it, I
don't recomand that you try to use it.
Also, Apple OpenCL is pretty buggy. Some basic openCL code have its
compiler crash. We have some workaround some of those cases, but not all of
them.
If you still want to try it, make sure to use the dev version of Theano.
On Tue, Aug 22, 2017 at 5:14 PM aditsanghvi94 notifications@github.com
wrote:
GPU: Radeon Pro 560, OS X
Theano is running on GPU via OpenCL, not CUDA. I ran the script here "
http://deeplearning.net/software/theano/tutorial/using_gpu.html"However, I am getting the following GPUArrayException error when I use
model.fit_generatorGpuArrayException Traceback (most recent call last)
in ()
76 epochs=epochs,
77 validation_data=validation_generator,
---> 78 validation_steps=nb_validation_samples // batch_size)
79
80/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/legacy/interfaces.pyc
in wrapper(args, *kwargs)
85 warnings.warn('Update your ' + object_name + 86 ' call to the Keras 2
API: ' + signature, stacklevel=2)
---> 87 return func(args, *kwargs)
88 wrapper._original_function = func
89 return wrapper/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/models.pyc in
fit_generator(self, generator, steps_per_epoch, epochs, verbose, callbacks,
validation_data, validation_steps, class_weight, max_queue_size, workers,
use_multiprocessing, initial_epoch)
1115 workers=workers,
1116 use_multiprocessing=use_multiprocessing,
-> 1117 initial_epoch=initial_epoch)
1118
1119 @interfaces.legacy_generator_methods_support/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/legacy/interfaces.pyc
in wrapper(args, *kwargs)
85 warnings.warn('Update your ' + object_name + 86 ' call to the Keras 2
API: ' + signature, stacklevel=2)
---> 87 return func(args, *kwargs)
88 wrapper._original_function = func
89 return wrapper/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/engine/training.pyc
in fit_generator(self, generator, steps_per_epoch, epochs, verbose,
callbacks, validation_data, validation_steps, class_weight, max_queue_size,
workers, use_multiprocessing, initial_epoch)
1723
1724 do_validation = bool(validation_data)
-> 1725 self._make_train_function()
1726 if do_validation:
1727 self._make_test_function()/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/engine/training.pyc
in _make_train_function(self)
942 updates=updates,
943 name='train_function',
--> 944 **self._function_kwargs)
945
946 def _make_test_function(self):/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/backend/theano_backend.pyc
in function(inputs, outputs, updates, *kwargs)
1204 msg = 'Invalid argument "%s" passed to K.function with Theano
backend' % key
1205 raise ValueError(msg)
-> 1206 return Function(inputs, outputs, updates=updates, *kwargs)
1207
1208/Users/sanghvia/anaconda/lib/python2.7/site-packages/keras/backend/theano_backend.pyc
in init(self, inputs, outputs, updates, name, *kwargs)
1190 on_unused_input='ignore',
1191 name=name,
-> 1192 *kwargs)
1193 self.name = name
1194/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/compile/function.pyc
in function(inputs, outputs, mode, updates, givens, no_default_updates,
accept_inplace, name, rebuild_strict, allow_input_downcast, profile,
on_unused_input)
324 on_unused_input=on_unused_input,
325 profile=profile,
--> 326 output_keys=output_keys)
327 # We need to add the flag check_aliased inputs if we have any mutable
or
328 # borrowed used defined inputs/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/compile/pfunc.pyc
in pfunc(params, outputs, mode, updates, givens, no_default_updates,
accept_inplace, name, rebuild_strict, allow_input_downcast, profile,
on_unused_input, output_keys)
484 accept_inplace=accept_inplace, name=name,
485 profile=profile, on_unused_input=on_unused_input,
--> 486 output_keys=output_keys)
487
488/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/compile/function_module.pyc
in orig_function(inputs, outputs, mode, accept_inplace, name, profile,
on_unused_input, output_keys)
1793 on_unused_input=on_unused_input,
1794 output_keys=output_keys).create(
-> 1795 defaults)
1796
1797 t2 = time.time()/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/compile/function_module.pyc
in create(self, input_storage, trustme, storage_map)
1659 theano.config.traceback.limit = theano.config.traceback.compile_limit
1660 _fn, _i, _o = self.linker.make_thunk(
-> 1661 input_storage=input_storage_lists, storage_map=storage_map)
1662 finally:
1663 theano.config.traceback.limit = limit_orig/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/link.pyc
in make_thunk(self, input_storage, output_storage, storage_map)
697 return self.make_all(input_storage=input_storage,
698 output_storage=output_storage,
--> 699 storage_map=storage_map)[:3]
700
701 def make_all(self, input_storage, output_storage):/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/vm.pyc in
make_all(self, profiler, input_storage, output_storage, storage_map)
1045 compute_map,
1046 no_recycling,
-> 1047 impl=impl))
1048 linker_make_thunk_time[node] = time.time() - thunk_start
1049 if not hasattr(thunks[-1], 'lazy'):/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/op.pyc in
make_thunk(self, node, storage_map, compute_map, no_recycling, impl)
933 try:
934 return self.make_c_thunk(node, storage_map, compute_map,
--> 935 no_recycling)
936 except (NotImplementedError, utils.MethodNotDefined):
937 # We requested the c code, so don't catch the error./Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/op.pyc in
make_c_thunk(self, node, storage_map, compute_map, no_recycling)
837 _logger.debug('Trying CLinker.make_thunk')
838 outputs = cl.make_thunk(input_storage=node_input_storage,
--> 839 output_storage=node_output_storage)
840 fill_storage, node_input_filters, node_output_filters = outputs
841/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in
make_thunk(self, input_storage, output_storage, storage_map, keep_lock)
1188 cthunk, in_storage, out_storage, error_storage = self.compile(
1189 input_storage, output_storage, storage_map,
-> 1190 keep_lock=keep_lock)
1191
1192 res = _CThunk(cthunk, init_tasks, tasks, error_storage)/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in
compile(self, input_storage, output_storage, storage_map, keep_lock)
1129 output_storage,
1130 storage_map,
-> 1131 keep_lock=keep_lock)
1132 return (thunk,
1133 [link.Container(input, storage) for input, storage in/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in
cthunk_factory(self, error_storage, in_storage, out_storage, storage_map,
keep_lock)
1584 node.op.prepare_node(node, storage_map, None, 'c')
1585 module = get_module_cache().module_from_key(
-> 1586 key=key, lnk=self, keep_lock=keep_lock)
1587
1588 vars = self.inputs + self.outputs + self.orphans/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cmodule.pyc
in module_from_key(self, key, lnk, keep_lock)
1120 return module
1121
-> 1122 src_code = lnk.get_src_code()
1123 # Is the source code already in the cache?
1124 module_hash = get_module_hash(src_code, key)/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in
get_src_code(self)
1460
1461 def get_src_code(self):
-> 1462 mod = self.get_dynamic_module()
1463 return mod.code()
1464/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in
get_dynamic_module(self)
1504 """
1505 if not hasattr(self, '_mod'):
-> 1506 self.code_gen()
1507
1508 mod = cmodule.DynamicModule()/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gof/cc.pyc in
code_gen(self)
780 try:
781 c_support_code_apply.append(op.c_support_code_apply(node,
--> 782 name))
783 except utils.MethodNotDefined:
784 pass/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gpuarray/basic_ops.pyc
in f(args)
450 def forward_string_meth(name):
451 def f(args):
--> 452 res = getattr(GpuKernelBase, name)(args)
453 try:
454 res = res + '\n' + getattr(COp, name)(args)/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gpuarray/basic_ops.pyc
in c_support_code_apply(self, node, name)
375 kernels = self.gpu_kernels(node, name)
376 ctx = self.get_params(node)
--> 377 bins = '\n'.join(self._generate_kernel_bin(k, ctx) for k in
kernels)
378 codes = '\n'.join(self._generate_kernel_code(k) for k in kernels)
379 return '\n'.join([bins, codes])/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gpuarray/basic_ops.pyc
in ((k,))
375 kernels = self.gpu_kernels(node, name)
376 ctx = self.get_params(node)
--> 377 bins = '\n'.join(self._generate_kernel_bin(k, ctx) for k in
kernels)
378 codes = '\n'.join(self._generate_kernel_code(k) for k in kernels)
379 return '\n'.join([bins, codes])/Users/sanghvia/anaconda/lib/python2.7/site-packages/theano/gpuarray/basic_ops.pyc
in _generate_kernel_bin(self, k, ctx)
317 def _generate_kernel_bin(self, k, ctx):
318 gk = gpuarray.GpuKernel(k.code, k.name, k.params, context=ctx,
--> 319 **k._get_py_flags())
320 bin = gk._binary
321 bcode = ','.join(hex(c) for c in iterbytes(bin))pygpu/gpuarray.pyx in pygpu.gpuarray.GpuKernel.cinit()
pygpu/gpuarray.pyx in pygpu.gpuarray.kernel_init()
GpuArrayException: ('The following error happened while compiling the
node', GpuCorrMM{valid, (1, 1), (1, 1)}(GpuContiguous.0, GpuContiguous.0),
'\n', u'Program build failure ::\n💯16: error: initializing 'float '
with an expression of type 'global float *' changes address space of
pointer\n DTYPE_i0 * data_col_ptr = data_col;\n ^~~\n:102:22: error:
initializing 'const float ' with an expression of type 'const __global
float ' changes address space of pointer\n const DTYPE_i0 * data_im_ptr =
data_im + data_im_offset;\n ^ ~~~~~~\n:109:17: warning:
comparison of unsigned expression >= 0 is always true\n (h_im >= 0 && w_im= 0 && h_im < height && w_im < width) ?\n
~~ ^ ~\n:109:30: warning:~ ^ ~\n0001\t#define
comparison of unsigned expression >= 0 is always true\n (h_im >= 0 && w_im
= 0 && h_im < height && w_im < width) ?\n ~
local_barrier() barrier(CLK_LOCAL_MEM_FENCE)\n0002\t#define WITHIN_KERNEL /
empty /\n0003\t#define KERNEL __kernel\n0004\t#define GLOBAL_MEM
__global\n0005\t#define LOCAL_MEM __local\n0006\t#define LOCAL_MEM_ARG
__local\n0007\t#ifndef NULL\n0008\t #define NULL
((void)0)\n0009\t#endif\n0010\t#define LID_0 get_local_id(0)\n0011\t#define
LID_1 get_local_id(1)\n0012\t#define LID_2 get_local_id(2)\n0013\t#define
LDIM_0 get_local_size(0)\n0014\t#define LDIM_1
get_local_size(1)\n0015\t#define LDIM_2 get_local_size(2)\n0016\t#define
GID_0 get_group_id(0)\n0017\t#define GID_1 get_group_id(1)\n0018\t#define
GID_2 get_group_id(2)\n0019\t#define GDIM_0
get_num_groups(0)\n0020\t#define GDIM_1 get_num_groups(1)\n0021\t#define
GDIM_2 get_num_groups(2)\n0022\t#define ga_bool uchar\n0023\t#define
ga_byte char\n0024\t#define ga_ubyte uchar\n0025\t#define ga_short
short\n0026\t#define ga_ushort ushort\n0027\t#define ga_int
int\n0028\t#define ga_uint uint\n0029\t#define ga_long long\n0030\t#define
ga_ulong ulong\n0031\t#define ga_float float\n0032\t#define ga_double
double\n0033\t#define ga_half half\n0034\t#define ga_size
ulong\n0035\t#define ga_ssize long\n0036\t#define load_half(p)
vload_half(0, p)\n0037\t#define store_half(p, v) vstore_half_rtn(v, 0,
p)\n0038\t#define GA_DECL_SHARED_PARAM(type, name) , __local type
name\n0039\t#define GA_DECL_SHARED_BODY(type, name)\n0040\t#define
GA_WARP_SIZE 64\n0041\t#define DTYPE_i0 ga_float\n0042\t#define DTYPE_i1
ga_float\n0043\t#define DTYPE_o0 ga_float\n0044\t\n0045\t\n0046\t// TODO
check kernel flags\n0047\t// This uses a lot of code from Caffe
(http://caffe.berkeleyvision.org/);\n0048\t//
http://caffe.berkeleyvision.org/);%5Cn0048%5Ct// sources are clearly
marked. Below we reproduce the original license of\n0049\t// the Caffe
software.\n0050\t/\n0051\tCopyright (c) 2014, The Regents of the University
of California (Regents)\n0052\tAll rights
reserved.\n0053\t\n0054\tRedistribution and use in source and binary forms,
with or without\n0055\tmodification, are permitted provided that the
following conditions are met:\n0056\t\n0057\t1. Redistributions of source
code must retain the above copyright notice, this\n0058\t list of
conditions and the following disclaimer.\n0059\t2. Redistributions in
binary form must reproduce the above copyright notice,\n0060\t this list of
conditions and the following disclaimer in the documentation\n0061\t and/or
other materials provided with the distribution.\n0062\t\n0063\tTHIS
SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
AND\n0064\tANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
TO, THE IMPLIED\n0065\tWARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
PARTICULAR PURPOSE ARE\n0066\tDISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
OWNER OR CONTRIBUTORS BE LIABLE FOR\n0067\tANY DIRECT, INDIRECT,
INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES\n0068\t(INCLUDING,
BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
SERVICES;\n0069\tLOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION)
HOWEVER CAUSED AND\n0070\tON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
STRICT LIABILITY, OR TORT\n0071\t(INCLUDING NEGLIGENCE OR OTHERWISE)
ARISING IN ANY WAY OUT OF THE USE OF THIS\n0072\tSOFTWARE, EVEN IF ADVISED
OF THE POSSIBILITY OF SUCH DAMAGE.\n0073\t/\n0074\t\n0075\t// (borrowed
from Caffe:
https://github.com/BVLC/caffe/blob/master/src/caffe/util/im2col.cu)\n0076\t//
https://github.com/BVLC/caffe/blob/master/src/caffe/util/im2col.cu)%5Cn0076%5Ct//
Kernels for fast unfold + copy\n0077\t// GPU kernel for the case of
dilation\n0078\tKERNEL void dilated_im2col_kernel(const ga_size n,\n0079\t
GLOBAL_MEM const DTYPE_i0 * data_im,\n0080\t const ga_size
data_im_offset,\n0081\t const ga_size height, const ga_size width,\n0082\t
const ga_size kernel_h, const ga_size kernelw,\n0083\t const ga_size
dilation_h, const ga_size dilationw,\n0084\t const ga_size pad_h, const
ga_size pad_w,\n0085\t const ga_size stride_h, const ga_size
stride_w,\n0086\t const ga_size height_col, const ga_size
width_col,\n0087\t GLOBAL_MEM DTYPE_i0 * data_col) {\n0088\t // grid stride
looping\n0089\t for (ga_size index = GID_0 * LDIM_0 + LID_0;\n0090\t index
< (n); index += LDIM_0 * GDIM_0) {\n0091\t const ga_size h_index = index /
width_col;\n0092\t const ga_size h_col = h_index % height_col;\n0093\t
const ga_size wcol = index % width_col;\n0094\t const ga_size c_im =
h_index / height_col;\n0095\t const ga_size c_col = c_im * kernel_h *
kernelw;\n0096\t const ga_size h_offset = h_col * stride_h -
pad_h;\n0097\t const ga_size woffset = wcol * stridew - padw;\n0098\t
DTYPE_i0 * data_col_ptr = data_col;\n0099\t data_col_ptr += (c_col *
height_col + h_col) * width_col + w_col;\n0100\t const DTYPE_i0 *
data_im_ptr = data_im + data_im_offset;\n0101\t data_im_ptr += (c_im *
height + h_offset) * width + woffset;\n0102\t for (ga_size i = 0; i <
kernel_h; ++i) {\n0103\t for (ga_size j = 0; j < kernelw; ++j) {\n0104\t
ga_size h_im = h_offset + i * dilation_h;\n0105\t ga_size w_im = woffset
- j * dilationw;\n0106\t *data_col_ptr =\n0107\t (h_im >= 0 && w_im >= 0
&& h_im < height && wim < width) ?\n0108\t data_im_ptr[i * dilation_h *
width + j * dilation*w] : 0;\n0109\t data_col_ptr += height_col *
width_col;\n0110\t }\n0111\t }\n0112\t
}\n0113\t}\n0114\t\n0115\t\n0116\t#undef DTYPE_i0\n0117\t#undef
DTYPE_i1\n0118\t#undef DTYPE_o0\n')I dug around a bit and found out that in feb 2017 someone said that theano
stopped openCL support. Is this still the case?—
You are receiving this because you are subscribed to this thread.
Reply to this email directly, view it on GitHub
https://github.com/Theano/Theano/issues/6343, or mute the thread
https://github.com/notifications/unsubscribe-auth/AALC-zWW0Zzthp9PRQTrSCrOEoyRxwPcks5sa0TGgaJpZM4O_Mim
.
nouiz
on 22 Aug 2017
As said here and elsewhere theano CL support is not complete.
Nonetheless, it is reported to work, at least for some tests. If it doesn't for you, is just because apple really, really hates standards.
mirh
on 15 Dec 2017
@mirh generally speaking, I couldn't agree more with bashing Apple for their standard aversion.
But your statement is ridiculous given that this is Apple that initiated the OpenCL standard and gave it to Khronos.
LifeIsStrange
on 5 Oct 2018
Related issues
Higgcz
·
6Comments
subzerofun
·
8Comments
9thDimension
·
3Comments
aloctavodia
·
8Comments
gammaguy
·
5Comments
Most helpful comment
We didn't stopped OpenCL support, we never finished it. We don't have time
to finish it, so unless you have a significant time to help us finish it, I
don't recomand that you try to use it.
Also, Apple OpenCL is pretty buggy. Some basic openCL code have its
compiler crash. We have some workaround some of those cases, but not all of
them.
If you still want to try it, make sure to use the dev version of Theano.
On Tue, Aug 22, 2017 at 5:14 PM aditsanghvi94 notifications@github.com
wrote: