diff --git a/hysop/backend/device/kernel_statistics.py b/hysop/backend/device/kernel_statistics.py index ed7484a5daf4094fa4b89fffeeb623c1defa329b..b37105488cced43f83640aca2d248c8a03631c17 100644 --- a/hysop/backend/device/kernel_statistics.py +++ b/hysop/backend/device/kernel_statistics.py @@ -7,7 +7,7 @@ class KernelStatistics(object): Execution statistics extracted from kernel events. """ - def __init__(self, min_, max_, total, nruns, data=None, **kwds): + def __init__(self, min_, max_, total, nruns, data=None, **kwds): """ Initialize KernelStatistics from nruns. Statistics should be given in nanoseconds. @@ -23,21 +23,21 @@ class KernelStatistics(object): self._total = total self._nruns = nruns self._data = None if (data is None) else tuple(data) - + def _get_min(self): - return self._min + return self._min def _get_max(self): - return self._max + return self._max def _get_mean(self): assert (self._nruns>0) return self._total / float(self._nruns) def _get_total(self): - return self._total + return self._total def _get_nruns(self): - return self._nruns + return self._nruns def _get_data(self): - return self._data - + return self._data + min = property(_get_min) max = property(_get_max) mean = property(_get_mean) @@ -77,7 +77,7 @@ class KernelStatistics(object): return self.cmp(self, other) >= 0 def __ne__(self, other): return self.cmp(self, other) != 0 - + def __iadd__(self, other): if (other.nruns == 0): return @@ -94,7 +94,7 @@ class KernelStatistics(object): self._total += other.total self._data += other.data return self - + def __str__(self): mini = float(self.min) * 1e-9 #ns maxi = float(self.max) * 1e-9 #ns diff --git a/hysop/backend/device/opencl/opencl_array.py b/hysop/backend/device/opencl/opencl_array.py index 42df5ec55a7474bf3d93c4a8cb4da618f547881c..41044931e5e26c538a067ef8cb88a283ca58debd 100644 --- a/hysop/backend/device/opencl/opencl_array.py +++ b/hysop/backend/device/opencl/opencl_array.py @@ -22,7 +22,7 @@ class OpenClArray(Array): handle: pyopencl.array.Array, implementation of this array kargs: arguments for base classes. """ - + if not isinstance(handle, clArray.Array): msg='Handle should be a pyopencl.array.Array but got a {}.' msg=msg.format(handle.__class__) @@ -31,28 +31,28 @@ class OpenClArray(Array): msg='Backend should be a OpenClArrayBackend but got a {}.' msg=msg.format(handle.__class__) raise ValueError(msg) - + if handle.dtype in [np.float16, np.longdouble, np.bool]: msg='{} unsupported yet for OpenCl arrays.'.format(handle.dtype) raise TypeError(msg) - + super(OpenClArray,self).__init__(handle=handle, backend=backend, **kargs) - - # at this time the opencl backend works only with the default_queue + + # at this time the opencl backend works only with the default_queue # so we enforce it. if (handle.queue is not None) and (handle.queue is not self.default_queue): msg = 'pyopencl.Array has been created with a non-default queue.' raise RuntimeError(msg) backend.check_queue(handle.queue) self.set_default_queue(self.default_queue) - + def as_symbolic_array(self, name, **kwds): """ Return a symbolic array variable that contain a reference to this array. """ from hysop.symbolic.array import OpenClSymbolicArray return OpenClSymbolicArray(memory_object=self, name=name, **kwds) - + def as_symbolic_buffer(self, name, **kwds): """ Return a symbolic buffer variable that contain a reference to this array. @@ -105,8 +105,8 @@ class OpenClArray(Array): return self._handle.nbytes def get_int_ptr(self): return self._handle.base_data.int_ptr + self.offset - - # array properties + + # array properties ndim = property(get_ndim) shape = property(get_shape, set_shape) offset = property(get_offset) @@ -121,7 +121,7 @@ class OpenClArray(Array): size = property(get_size) nbytes = property(get_nbytes) int_ptr = property(get_int_ptr) - + def get_base_data(self): return self._handle.base_data def get_offset(self): @@ -134,10 +134,10 @@ class OpenClArray(Array): Equivalent C type corresponding to the numpy.dtype. """ return clTools.dtype_to_ctype(self.dtype) - - + + def get(self, handle=False, - queue=None, ary=None, async=False): + queue=None, ary=None, synchronize=True): """ Returns a HostArray, view or copy of this array. """ @@ -145,16 +145,16 @@ class OpenClArray(Array): if self.size==0: return None elif self.flags.forc: - host_array = self._call('get', queue=queue, ary=ary, async=async) + host_array = self._call('get', queue=queue, ary=ary, synchronize=synchronize) else: from hysop.backend.device.opencl.opencl_copy_kernel_launchers import \ OpenClCopyBufferRectLauncher - assert not async + assert synchronize if (ary is not None): host_array = ary else: host_array = self.backend.host_array_backend.empty_like(self) - kl = OpenClCopyBufferRectLauncher.from_slices(varname='buffer', src=self, + kl = OpenClCopyBufferRectLauncher.from_slices(varname='buffer', src=self, dst=host_array) evt = kl(queue=queue) evt.wait() @@ -167,19 +167,19 @@ class OpenClArray(Array): if host_array.ndim == 1: return host_array._handle[0] return host_array - + # event managment def events(self): """ - A list of pyopencl.Event instances that the current content of this array depends on. - User code may read, but should never modify this list directly. + A list of pyopencl.Event instances that the current content of this array depends on. + User code may read, but should never modify this list directly. To update this list, instead use the following methods. """ return self.handle.events def add_event(self, evt): """ - Add evt to events. If events is too long, this method may implicitly wait + Add evt to events. If events is too long, this method may implicitly wait for a subset of events and clear them from the list. """ self._call('add_event', evt=evt) @@ -205,7 +205,7 @@ class OpenClArray(Array): """ Sets the default queue for this array. """ - # at this time the opencl backend works only with the default_queue + # at this time the opencl backend works only with the default_queue # so we enforce it. if (queue is not self.default_queue): msg='Default queue override has been disabled for non-default queues.' @@ -222,7 +222,7 @@ class OpenClArray(Array): Get the default queue for this array. """ return self._handle.queue or self.backend.default_queue - + context = property(get_context) device = property(get_device) default_queue = property(get_default_queue, set_default_queue) @@ -234,25 +234,25 @@ class OpenClArray(Array): queue = self.backend.check_queue(queue) yield self._call('with_queue', queue=queue) - + ## Array specific methods def view(self, dtype=None): """ - Returns view of array with the same data. If dtype is different from current dtype, + Returns view of array with the same data. If dtype is different from current dtype, the actual bytes of memory will be reinterpreted. """ return self._call('view', dtype=dtype) - + def reshape(self, shape, order=default_order): """ - Returns view of array with the same data. If dtype is different from current dtype, + Returns view of array with the same data. If dtype is different from current dtype, the actual bytes of memory will be reinterpreted. """ shape = tuple(int(i) for i in shape) return self._call('reshape', *shape, order=order) def astype(self, dtype, queue=None, - order=MemoryOrdering.SAME_ORDER, + order=MemoryOrdering.SAME_ORDER, casting='unsafe', subok=True, copy=True): """ Copy of the array, cast to a specified type. @@ -263,119 +263,119 @@ class OpenClArray(Array): self._unsupported_argument('astype', 'copy', copy, True) queue = self.backend.check_queue(queue) return self._call('astype', dtype=dtype, queue=queue) - + ## Cached kernels for efficiency - def min(self, axis=None, out=None, - queue=None, async=False, **kwds): + def min(self, axis=None, out=None, + queue=None, synchronize=True, **kwds): """ Return the minimum along a given axis. On the first call, a kernel launcher is built for efficiency. """ if (axis is None) and (out is None): if not hasattr(self, '_OpenClArray__min_launcher'): - self.__min_launcher = self.backend.amin(a=self, axis=axis, out=out, - build_kernel_launcher=True, queue=queue, async=True, + self.__min_launcher = self.backend.amin(a=self, axis=axis, out=out, + build_kernel_launcher=True, queue=queue, synchronize=False, **kwds) - evt = self.__min_launcher(queue=queue, async=True) + evt = self.__min_launcher(queue=queue, synchronize=False) out = self.__min_launcher.out - if async: - return evt - else: + if synchronize: evt.wait() return out.copy() + else: + return evt else: super(OpenClArray, self).min(self, axis=axis, out=out, - queue=queue, async=async, **kwds) + queue=queue, synchronize=synchronize, **kwds) - def max(self, axis=None, out=None, - queue=None, async=False, **kwds): + def max(self, axis=None, out=None, + queue=None, synchronize=True, **kwds): """ Return the maximum along a given axis. On the first call, a kernel launcher is built for efficiency. """ if (axis is None) and (out is None): if not hasattr(self, '_OpenClArray__max_launcher'): - self.__max_launcher = self.backend.amax(a=self, axis=axis, out=out, - build_kernel_launcher=True, queue=queue, async=True, + self.__max_launcher = self.backend.amax(a=self, axis=axis, out=out, + build_kernel_launcher=True, queue=queue, synchronize=False, **kwds) - evt = self.__max_launcher(queue=queue, async=True) + evt = self.__max_launcher(queue=queue, synchronize=False) out = self.__max_launcher.out - if async: - return evt - else: + if synchronize: evt.wait() return out.copy() + else: + return evt else: super(OpenClArray, self).max(self, axis=axis, out=out, - queue=queue, async=async, **kwds) - - def nanmin(self, axis=None, out=None, - queue=None, async=False, **kwds): + queue=queue, synchronize=synchronize, **kwds) + + def nanmin(self, axis=None, out=None, + queue=None, synchronize=True, **kwds): """ Return the minimum along a given axis. On the first call, a kernel launcher is built for efficiency. """ if (axis is None) and (out is None): if not hasattr(self, '_OpenClArray__nanmin_launcher'): - self.__nanmin_launcher = self.backend.nanmin(a=self, axis=axis, out=out, - build_kernel_launcher=True, queue=queue, async=True, + self.__nanmin_launcher = self.backend.nanmin(a=self, axis=axis, out=out, + build_kernel_launcher=True, queue=queue, synchronize=False, **kwds) - evt = self.__nanmin_launcher(queue=queue, async=True) + evt = self.__nanmin_launcher(queue=queue, synchronize=False) out = self.__nanmin_launcher.out - if async: - return evt - else: + if synchronize: evt.wait() return out.copy() + else: + return evt else: super(OpenClArray, self).nanmin(self, axis=axis, out=out, - queue=queue, async=async, **kwds) + queue=queue, synchronize=synchronize, **kwds) - def nanmax(self, axis=None, out=None, - queue=None, async=False, **kwds): + def nanmax(self, axis=None, out=None, + queue=None, synchronize=True, **kwds): """ Return the maximum along a given axis. On the first call, a kernel launcher is built for efficiency. """ if (axis is None) and (out is None): if not hasattr(self, '_OpenClArray__nanmax_launcher'): - self.__nanmax_launcher = self.backend.nanmax(a=self, axis=axis, out=out, - build_kernel_launcher=True, queue=queue, async=True, + self.__nanmax_launcher = self.backend.nanmax(a=self, axis=axis, out=out, + build_kernel_launcher=True, queue=queue, synchronize=False, **kwds) - evt = self.__nanmax_launcher(queue=queue, async=True) + evt = self.__nanmax_launcher(queue=queue, synchronize=False) out = self.__nanmax_launcher.out - if async: - return evt - else: + if synchronize: evt.wait() return out.copy() + else: + return evt else: super(OpenClArray, self).nanmax(self, axis=axis, out=out, - queue=queue, async=async, **kwds) - - def sum(self, axis=None, out=None, - queue=None, async=False, **kwds): + queue=queue, synchronize=synchronize, **kwds) + + def sum(self, axis=None, out=None, + queue=None, synchronize=True, **kwds): """ Return the sum along a given axis. On the first call, a kernel launcher is built for efficiency. """ if (axis is None) and (out is None): if not hasattr(self, '_OpenClArray__sum_launcher'): - self.__sum_launcher = self.backend.sum(a=self, axis=axis, out=out, - build_kernel_launcher=True, queue=queue, async=True, + self.__sum_launcher = self.backend.sum(a=self, axis=axis, out=out, + build_kernel_launcher=True, queue=queue, synchronize=False, **kwds) - evt = self.__sum_launcher(queue=queue, async=True) + evt = self.__sum_launcher(queue=queue, synchronize=False) out = self.__sum_launcher.out - if async: - return evt - else: + if synchronize: evt.wait() return out.copy() + else: + return evt else: super(OpenClArray, self).sum(self, axis=axis, out=out, - queue=queue, async=async, **kwds) - + queue=queue, synchronize=synchronize, **kwds) + def setitem(self, subscript, value, queue=None): queue = first_not_None(queue, self.default_queue) if np.isscalar(value): @@ -384,12 +384,12 @@ class OpenClArray(Array): a.fill(value=value, queue=queue) else: try: - self.handle.setitem(subscript=subscript, value=value, + self.handle.setitem(subscript=subscript, value=value, queue=queue) except: from hysop.backend.device.opencl.opencl_copy_kernel_launchers import \ OpenClCopyBufferRectLauncher - kl = OpenClCopyBufferRectLauncher.from_slices(varname='buffer', src=value, + kl = OpenClCopyBufferRectLauncher.from_slices(varname='buffer', src=value, dst=self, dst_slices=subscript) evt = kl(queue=queue) evt.wait() @@ -398,7 +398,7 @@ class OpenClArray(Array): if any( (s==0) for s in self[subscript].shape ): return self.setitem(subscript=subscript, value=value, **kwds) - + def __str__(self): return str(self.get()) def __repr__(self): diff --git a/hysop/backend/device/opencl/opencl_array_backend.py b/hysop/backend/device/opencl/opencl_array_backend.py index bd4e9c3bda3905f3b77a73ae3a5752acd43e36d6..cdaaf11126e343046811031706d854b4cb23b06d 100644 --- a/hysop/backend/device/opencl/opencl_array_backend.py +++ b/hysop/backend/device/opencl/opencl_array_backend.py @@ -469,7 +469,7 @@ class OpenClArrayBackend(ArrayBackend): return super(OpenClArrayBackend,self)._arg(arg) - def copyto(self, dst, src, queue=None, async=False, **kargs): + def copyto(self, dst, src, queue=None, synchronize=True, **kargs): """ src is a OpenClArray dst can be everything @@ -493,14 +493,14 @@ class OpenClArrayBackend(ArrayBackend): kl = OpenClCopyBufferRectLauncher.from_slices('buffer', src=src, dst=dst) evt = kl(queue=queue) - if async: - return evt - else: + if synchronize: evt.wait() + else: + return evt elif isinstance(dst, Array): - return src.handle.get(queue=queue, ary=dst.handle, async=async) + return src.handle.get(queue=queue, ary=dst.handle, synchronize=synchronize) elif isinstance(dst, np.ndarray): - return src.handle.get(queue=queue, ary=dst, async=async) + return src.handle.get(queue=queue, ary=dst, synchronize=synchronize) else: msg = 'Unknown type to copy to ({}) for array of type {}.' msg = msg.format(dst.__class__.__name__, src.__class__.__name__) @@ -748,7 +748,7 @@ class OpenClArrayBackend(ArrayBackend): filter_expr=None, dtype=None, infer_dtype=False, queue=None, infer_queue=False, - async=False, convert_inputs=None, + synchronize=True, convert_inputs=None, allocator=None, alloc_shapes=None, alloc_dtypes=None, Kernel=_ElementwiseKernel, kernel_build_kwargs=None, kernel_call_kwargs=None, @@ -848,10 +848,10 @@ class OpenClArrayBackend(ArrayBackend): Call and events: wait_for specify a list a event to wait prior applying the kernel. - If async is set to False, this is a blocking opencl call and this functions returns + If syncrhonize is set, this is a blocking opencl call and this functions returns output arguments okargs as a tuple unless there is only one output. - If async is set to True, this is a non blocking call and an event is returned in + If syncrhonize not set, this is a non blocking call and an event is returned in addition to the outputs as last argument. The returned event may not be used for profiling purposes, because it only @@ -1020,7 +1020,7 @@ class OpenClArrayBackend(ArrayBackend): okargs = self._return(okargs) - if not async: + if synchronize: evt.wait() if len(okargs)==1: return okargs[0] @@ -1039,7 +1039,7 @@ class OpenClArrayBackend(ArrayBackend): queue=None, infer_queue=True, allocator=None, alloc_shapes=None, alloc_dtypes=None, - async=False, wait_for=None, + synchronize=True, wait_for=None, convert_inputs=None, build_kernel_launcher=False, name='nary_op', options=[], preamble=''): @@ -1060,7 +1060,7 @@ class OpenClArrayBackend(ArrayBackend): <=> y0[i] = ElementwiseKernel(const x0[i],...,const xn[i], const p0,..., const pk) - If async is set to True, evt is returned as last argument. + If synchronize is set to false, evt is returned as last argument. If out and dtype is None it is set to xi.dtype where xi is the first OpenClArray options are opencl kernel build options. @@ -1084,7 +1084,7 @@ class OpenClArrayBackend(ArrayBackend): filter_expr=filter_expr, dtype=dtype, infer_dtype=infer_dtype, queue=queue, infer_queue=infer_queue, - async=async, + synchronize=synchronize, allocator=allocator, convert_inputs=convert_inputs, alloc_shapes=alloc_shapes, alloc_dtypes=alloc_dtypes, Kernel=Kernel, @@ -1100,7 +1100,7 @@ class OpenClArrayBackend(ArrayBackend): convert_inputs=None, allocator=None, alloc_shapes=None, alloc_dtypes=None, - async=False, wait_for=None, + synchronize=True, wait_for=None, build_kernel_launcher=False, name='unary_op', options=[], preamble=''): @@ -1111,7 +1111,7 @@ class OpenClArrayBackend(ArrayBackend): queue=queue, infer_queue=infer_queue, allocator=allocator, convert_inputs=convert_inputs, alloc_shapes=alloc_shapes, alloc_dtypes=alloc_dtypes, - async=async, wait_for=wait_for, + synchronize=synchronize, wait_for=wait_for, build_kernel_launcher=build_kernel_launcher, name=name, options=options, preamble=preamble) @@ -1123,7 +1123,7 @@ class OpenClArrayBackend(ArrayBackend): convert_inputs=None, allocator=None, alloc_shapes=None, alloc_dtypes=None, - async=False, wait_for=None, + synchronize=True, wait_for=None, build_kernel_launcher=False, name='binary_op', options=[], preamble=''): @@ -1134,7 +1134,7 @@ class OpenClArrayBackend(ArrayBackend): queue=queue, infer_queue=infer_queue, allocator=allocator, convert_inputs=convert_inputs, alloc_shapes=alloc_shapes, alloc_dtypes=alloc_dtypes, - async=async, wait_for=wait_for, + synchronize=synchronize, wait_for=wait_for, build_kernel_launcher=build_kernel_launcher, name=name, options=options, preamble=preamble) @@ -1147,7 +1147,7 @@ class OpenClArrayBackend(ArrayBackend): dtype=None, infer_dtype=True, queue=None, infer_queue=True, allocator=None, - async=False, wait_for=None, + synchronize=True, wait_for=None, build_kernel_launcher=False, name='reduction', options=[], preamble=''): """ @@ -1207,7 +1207,7 @@ class OpenClArrayBackend(ArrayBackend): filter_expr=filter_expr, dtype=dtype, infer_dtype=infer_dtype, queue=queue, infer_queue=infer_queue, - async=async, + synchronize=synchronize, convert_inputs=convert_inputs, alloc_dtypes=alloc_dtypes, allocator=allocator, alloc_shapes=((1,),), Kernel=Kernel, @@ -1224,7 +1224,7 @@ class OpenClArrayBackend(ArrayBackend): queue=None, infer_queue=True, allocator=None, convert_inputs=None, - async=False, wait_for=None, + synchronize=True, wait_for=None, build_kernel_launcher=False, name='nan_reduction', options=[], preamble=''): """ @@ -1248,7 +1248,7 @@ class OpenClArrayBackend(ArrayBackend): axis=axis, out=out, extra_kargs=extra_kargs, extra_arguments=extra_arguments, input_arguments=input_arguments, output_arguments=output_arguments, dtype=dtype, infer_dtype=infer_dtype, queue=queue, infer_queue=infer_queue, - allocator=allocator, async=async, wait_for=wait_for, + allocator=allocator, synchronize=synchronize, wait_for=wait_for, build_kernel_launcher=build_kernel_launcher, name=name, options=options, preamble=preamble) @@ -1262,7 +1262,7 @@ class OpenClArrayBackend(ArrayBackend): queue=None, infer_queue=True, allocator=None, convert_inputs=None, - async=False, wait_for=None, + synchronize=True, wait_for=None, build_kernel_launcher=False, name='generic_scan', options=[], preamble=''): """ @@ -1325,7 +1325,7 @@ class OpenClArrayBackend(ArrayBackend): filter_expr=filter_expr, dtype=dtype, infer_dtype=infer_dtype, queue=queue, infer_queue=infer_queue, - async=async, + synchronize=synchronize, allocator=allocator, alloc_shapes=None, convert_inputs=convert_inputs, Kernel=Kernel, @@ -1342,7 +1342,7 @@ class OpenClArrayBackend(ArrayBackend): queue=None, infer_queue=True, allocator=None, convert_inputs=None, - async=False, wait_for=None, + synchronize=True, wait_for=None, build_kernel_launcher=False, name='inclusive_scan', options=[], preamble=''): @@ -1357,7 +1357,7 @@ class OpenClArrayBackend(ArrayBackend): dtype=dtype, infer_dtype=infer_dtype, queue=queue, infer_queue=infer_queue, allocator=allocator, convert_inputs=convert_inputs, - async=async, wait_for=wait_for, + synchronize=synchronize, wait_for=wait_for, build_kernel_launcher=build_kernel_launcher, name=name, options=options, preamble=preamble) @@ -1369,7 +1369,7 @@ class OpenClArrayBackend(ArrayBackend): dtype=None, infer_dtype=True, queue=None, infer_queue=True, allocator=None, - async=False, wait_for=None, + synchronize=True, wait_for=None, build_kernel_launcher=False, name='exclusive_scan', options=[], preamble=''): @@ -1384,7 +1384,7 @@ class OpenClArrayBackend(ArrayBackend): dtype=dtype, infer_dtype=infer_dtype, queue=queue, infer_queue=infer_queue, allocator=allocator, - async=async, wait_for=wait_for, + synchronize=synchronize, wait_for=wait_for, build_kernel_launcher=build_kernel_launcher, name=name, options=options, preamble=preamble) @@ -1396,7 +1396,7 @@ class OpenClArrayBackend(ArrayBackend): dtype=None, infer_dtype=True, queue=None, infer_queue=True, allocator=None, - async=False, wait_for=None, + synchronize=True, wait_for=None, build_kernel_launcher=False, name='inclusive_nanscan', options=[], preamble=''): @@ -1419,7 +1419,7 @@ class OpenClArrayBackend(ArrayBackend): dtype=dtype, infer_dtype=infer_dtype, queue=queue, infer_queue=infer_queue, allocator=allocator, - async=async, wait_for=wait_for, + synchronize=synchronize, wait_for=wait_for, build_kernel_launcher=build_kernel_launcher, name=name, options=options, preamble=preamble) @@ -1431,7 +1431,7 @@ class OpenClArrayBackend(ArrayBackend): dtype=None, infer_dtype=True, queue=None, infer_queue=True, allocator=None, - async=False, wait_for=None, + synchronize=True, wait_for=None, build_kernel_launcher=False, name='exclusive_nanscan', options=[], preamble=''): @@ -1454,7 +1454,7 @@ class OpenClArrayBackend(ArrayBackend): dtype=dtype, infer_dtype=infer_dtype, queue=queue, infer_queue=infer_queue, allocator=allocator, - async=async, wait_for=wait_for, + synchronize=synchronize, wait_for=wait_for, build_kernel_launcher=build_kernel_launcher, name=name, options=options, preamble=preamble) @@ -1527,7 +1527,7 @@ class OpenClArrayBackend(ArrayBackend): array = self.wrap(handle) return array - def asarray(self, a, queue=None, async=False, + def asarray(self, a, queue=None, synchronize=True, dtype=None, order=default_order, array_queue=QueuePolicy.SAME_AS_TRANSFER): """ Convert the input to an OpenClArray. @@ -1560,7 +1560,7 @@ class OpenClArrayBackend(ArrayBackend): a = a.handle array = self._call(clArray.to_device, queue=queue, ary=a, - async=async, array_queue=array_queue) + synchronize=synchronize, array_queue=array_queue) else: msg='Unknown type to convert from {}.' msg=msg.format(acls) diff --git a/hysop/backend/device/opencl/opencl_env.py b/hysop/backend/device/opencl/opencl_env.py index 03617c15ed6c2895439a77382582016320825d9c..b3b6198ca48dbe8462d4d02d02f628e8569a1c88 100644 --- a/hysop/backend/device/opencl/opencl_env.py +++ b/hysop/backend/device/opencl/opencl_env.py @@ -510,7 +510,7 @@ Dumped OpenCL Kernel '{}' # --- Build kernel --- try: build = prg.build(build_options) - except Exception, e: + except Exception as e: print("Build files : ") for sf in file_list: print(" - ", sf) diff --git a/hysop/backend/device/opencl/operator/enstrophy.py b/hysop/backend/device/opencl/operator/enstrophy.py index c587f106276472f552b2debb08e9808a37f74e5c..c067736e2a6ee7865cd5b8c1e66660f6bdf9ad74 100644 --- a/hysop/backend/device/opencl/operator/enstrophy.py +++ b/hysop/backend/device/opencl/operator/enstrophy.py @@ -48,7 +48,7 @@ class OpenClEnstrophy(EnstrophyBase, OpenClSymbolic): super(OpenClEnstrophy, self).setup(work) (self.WdotW_kernel, self.WdotW_update_parameters) = self.symbolic_kernels['WdotW'] self.sum_kernel = self.dWdotW.backend.sum(a=self.dWdotW.sdata, - build_kernel_launcher=True, async=True) + build_kernel_launcher=True, synchronize=False) @op_apply def apply(self, **kwds): diff --git a/hysop/backend/device/opencl/operator/integrate.py b/hysop/backend/device/opencl/operator/integrate.py index 166d126e468f1de11c440e59d0b6ffa53062e5a6..5682a5b62a2b85945ac2ce018e070373ee99de66 100644 --- a/hysop/backend/device/opencl/operator/integrate.py +++ b/hysop/backend/device/opencl/operator/integrate.py @@ -24,7 +24,7 @@ class OpenClIntegrate(IntegrateBase, OpenClOperator): if self.expr is None: self.sum_kernels = tuple( self.dF.backend.sum(a=self.dF.data[i], - build_kernel_launcher=True, async=True) + build_kernel_launcher=True, synchronize=False) for i in range(self.dF.nb_components)) else: from hysop.backend.device.codegen.base.variables import dtype_to_ctype diff --git a/hysop/fields/tests/test_fields.py b/hysop/fields/tests/test_fields.py index 9e7aead6e84b9f4b16796b3664c7151ae9a61c48..a844654a5173425dcbe46e4399269df2cf5f409f 100644 --- a/hysop/fields/tests/test_fields.py +++ b/hysop/fields/tests/test_fields.py @@ -63,8 +63,8 @@ def test_field(): assert F0.domain is domain assert F0.dim is domain.dim - assert F0.name is 'F0' - assert F0.pretty_name is 'F0' + assert F0.name == 'F0' + assert F0.pretty_name == 'F0' assert F0 is not F1 assert F0 == F0 assert F0 != F1 @@ -73,8 +73,8 @@ def test_field(): assert hash(F0) != hash(F2) assert str(F0) == F0.long_description() assert len(F0.short_description()) < len(F0.long_description()) - assert F1.name is 'F1' - assert F1.pretty_name is 'test' + assert F1.name == 'F1' + assert F1.pretty_name == 'test' assert F0.nb_components == 1 assert len(F0.fields) == 1 assert F0.fields[0] is F0 diff --git a/hysop/numerics/fft/_mkl_fft.py b/hysop/numerics/fft/_mkl_fft.py index 91c479af0fc4f10251ca7d742e7a33639f1a93fb..1d1758369487cb39dfbc2d9acdc536ef76f1968d 100644 --- a/hysop/numerics/fft/_mkl_fft.py +++ b/hysop/numerics/fft/_mkl_fft.py @@ -60,7 +60,7 @@ def setup_transform(x, axis, transform, inverse, type): dout = dtype else: raise NotImplementedError - elif (transform is 'dst'): + elif (transform == 'dst'): ctype = float_to_complex_dtype(x.dtype) if (type==1): sin = mk_shape(shape, axis, 2*N+2) diff --git a/hysop/numerics/fft/gpyfft_fft.py b/hysop/numerics/fft/gpyfft_fft.py index b3a23b41e1ba59bab653cec6e9a9a25d127bb4f1..e49ffecd98856225a08e99815665404f9622a3e3 100644 --- a/hysop/numerics/fft/gpyfft_fft.py +++ b/hysop/numerics/fft/gpyfft_fft.py @@ -279,7 +279,7 @@ class GpyFFTPlan(OpenClFFTPlanI): plan.batch_size = t_batchsize_in plan.precision = t_precision plan.layouts = (layout_in, layout_out) - if (scaling is 'DEFAULT'): + if (scaling == 'DEFAULT'): pass elif (scaling is not None): plan.scale_forward = scale @@ -368,7 +368,7 @@ Post callback source code: self.pre_callback_src, self.post_callback_src) print(msg) - if (scaling is 'DEFAULT'): + if (scaling == 'DEFAULT'): pass elif (scaling is not None): plan.scale_forward = scale diff --git a/hysop/tools/decorators.py b/hysop/tools/decorators.py index 2bbe977f20f6d7eb29605a30ff6d6417467ab266..6b57a914438289ac626d5f7516beb26c6f6cb4be 100644 --- a/hysop/tools/decorators.py +++ b/hysop/tools/decorators.py @@ -73,7 +73,7 @@ def debug(f): if 'name' in kw: print('with name {}.'.format(kw['name'])) - if f.__name__ is '__new__': + if f.__name__ == '__new__': fullclassname = args[0].__mro__[0].__module__ + '.' fullclassname += args[0].__mro__[0].__name__ print('=> {}'.format(fullclassname)) diff --git a/hysop/topology/cartesian_topology.py b/hysop/topology/cartesian_topology.py index b7443e3f9b482dbdf70c4418477fbd5f634d5efa..c3105ef62fb3d0cf079b26f42bd372f128cbec86 100644 --- a/hysop/topology/cartesian_topology.py +++ b/hysop/topology/cartesian_topology.py @@ -79,9 +79,9 @@ class CartesianTopologyState(TopologyState): def _set_memory_order(self, memory_order): """Set the current memory_order as a tuple of hysop.constants.memory_order.""" memory_order = first_not_None(memory_order, MemoryOrdering.C_CONTIGUOUS) - if (memory_order is 'c'): + if (memory_order == 'c'): memory_order = MemoryOrdering.C_CONTIGUOUS - elif (memory_order is 'f'): + elif (memory_order == 'f'): memory_order = MemoryOrdering.F_CONTIGUOUS assert memory_order in (MemoryOrdering.C_CONTIGUOUS, MemoryOrdering.F_CONTIGUOUS), memory_order