diff --git a/hysop/backend/device/opencl/opencl_array_backend.py b/hysop/backend/device/opencl/opencl_array_backend.py index 15cbbe2ed0592810757f6c93aa3326868d419dbe..4bff1cfa93a59ae20c7d881a4b4950f981863a98 100644 --- a/hysop/backend/device/opencl/opencl_array_backend.py +++ b/hysop/backend/device/opencl/opencl_array_backend.py @@ -41,13 +41,12 @@ class _ElementwiseKernel(object): """ def __init__(self, context, arguments, operation, name='elementwise', preamble='', options=[], - **kargs): + **kwds): self.kernel = clElementwise.ElementwiseKernel(context=context, arguments=arguments, operation=operation, name=name, preamble=preamble, options=options) - def __call__(self, wait_for, args, **kargs): - return self.kernel.__call__(wait_for=wait_for, *args) - + def __call__(self, wait_for, args, **kwds): + return self.kernel.__call__(wait_for=wait_for, queue=kwds.get('queue',None), *args) def to_kernel_launcher(self, name, wait_for, queue, args, **kwds): """Build an _OpenClElementWiseKernelLauncher from self.""" @@ -889,6 +888,9 @@ class OpenClArrayBackend(ArrayBackend): if (shape is None): msg='No OpenClArray argument was found in ikargs and okargs.' raise ValueError(msg) + if (queue is None): + msg='Queue has not been specified.' + raise ValueError(msg) alloc_dtypes = tuple([match_dtype(dtype,alloc_dtypes[i]) \ for i in xrange(len(okargs))]) diff --git a/hysop/core/arrays/tests/test_array.py b/hysop/core/arrays/tests/test_array.py index c470124a2e1eec0f21dbe2c37ac332e0cb900ba0..39eaf54d0ee5cc6f0a7b95fc1f21b252b1617487 100644 --- a/hysop/core/arrays/tests/test_array.py +++ b/hysop/core/arrays/tests/test_array.py @@ -465,11 +465,9 @@ class TestArray(object): class TestContext(object): def __init__(self, opname, input_constraints, variables): - self.opname = opname - - # if there is a specific constrain we copy everything + # if there is a specific constraint we copy everything dtypes = {} if opname in input_constraints: for vname,vargs in variables.iteritems(): @@ -526,9 +524,13 @@ class TestArray(object): r1 = r1.astype(np.float16) if r0.dtype==np.bool: r1 = r1.astype(np.bool) - - l2 = np.sqrt(np.nansum((r0-r1)*np.conj(r0-r1)) / r0.size) - linf = np.nanmax(np.abs(r0-r1)) + + if (r0.dtype == np.bool) and (r1.dtype == np.bool): + l2 = np.sqrt(np.nansum(r0^r1)) / r0.size + linf = np.nanmax(r0^r1) + else: + l2 = np.sqrt(np.nansum((r0-r1)*np.conj(r0-r1)) / r0.size) + linf = np.nanmax(np.abs(r0-r1)) msg1='(l2={}, linf={}).' msg1=msg1.format(l2, linf) @@ -789,17 +791,17 @@ class TestArray(object): for cl_env in iter_clenv(): print print 'TESTING OPENCL PLATFORM {}'.format(cl_env.platform.name) - allocator = OpenClImmediateAllocator(cl_env.default_queue) - backend = OpenClArrayBackend(cl_env=cl_env, allocator=allocator) + allocator = OpenClImmediateAllocator(queue=cl_env.default_queue) + backend = OpenClArrayBackend(cl_env=cl_env, allocator=allocator) self._test_backend(backend) @opencl_failed def test_opencl_array_backend_pool(self): from hysop.backend.device.opencl.opencl_allocator import OpenClImmediateAllocator for cl_env in iter_clenv(): - allocator = OpenClImmediateAllocator(cl_env.default_queue)\ + allocator = OpenClImmediateAllocator(queue=cl_env.default_queue)\ .memory_pool(name=cl_env.device.name) - backend = OpenClArrayBackend(cl_env=cl_env, allocator=allocator) + backend = OpenClArrayBackend(cl_env=cl_env, allocator=allocator) self._test_backend(backend)