diff options
Diffstat (limited to 'silx/opencl/test')
-rw-r--r-- | silx/opencl/test/test_addition.py | 28 | ||||
-rw-r--r-- | silx/opencl/test/test_backprojection.py | 3 | ||||
-rw-r--r-- | silx/opencl/test/test_convolution.py | 99 |
3 files changed, 70 insertions, 60 deletions
diff --git a/silx/opencl/test/test_addition.py b/silx/opencl/test/test_addition.py index 49cc0b4..19dfdf0 100644 --- a/silx/opencl/test/test_addition.py +++ b/silx/opencl/test/test_addition.py @@ -29,19 +29,17 @@ Simple test of an addition """ -from __future__ import division, print_function - __authors__ = ["Henri Payno, Jérôme Kieffer"] __contact__ = "jerome.kieffer@esrf.eu" __license__ = "MIT" __copyright__ = "2013 European Synchrotron Radiation Facility, Grenoble, France" -__date__ = "01/08/2019" +__date__ = "30/11/2020" import logging import numpy import unittest -from ..common import ocl, _measure_workgroup_size +from ..common import ocl, _measure_workgroup_size, query_kernel_info if ocl: import pyopencl import pyopencl.array @@ -116,7 +114,7 @@ class TestAddition(unittest.TestCase): @unittest.skipUnless(ocl, "pyopencl is missing") def test_measurement(self): """ - tests that all devices are working properly ... + tests that all devices are working properly ... lengthy and error prone """ for platform in ocl.platforms: for did, device in enumerate(platform.devices): @@ -124,11 +122,31 @@ class TestAddition(unittest.TestCase): self.assertEqual(meas, device.max_work_group_size, "Workgroup size for %s/%s: %s == %s" % (platform, device, meas, device.max_work_group_size)) + @unittest.skipUnless(ocl, "pyopencl is missing") + def test_query(self): + """ + tests that all devices are working properly ... lengthy and error prone + """ + for what in ("COMPILE_WORK_GROUP_SIZE", + "LOCAL_MEM_SIZE", + "PREFERRED_WORK_GROUP_SIZE_MULTIPLE", + "PRIVATE_MEM_SIZE", + "WORK_GROUP_SIZE"): + logger.info("%s: %s", what, query_kernel_info(program=self.program, kernel="addition", what=what)) + + # Not all ICD work properly .... + #self.assertEqual(3, len(query_kernel_info(program=self.program, kernel="addition", what="COMPILE_WORK_GROUP_SIZE")), "3D kernel") + + min_wg = query_kernel_info(program=self.program, kernel="addition", what="PREFERRED_WORK_GROUP_SIZE_MULTIPLE") + max_wg = query_kernel_info(program=self.program, kernel="addition", what="WORK_GROUP_SIZE") + self.assertEqual(max_wg % min_wg, 0, msg="max_wg is a multiple of min_wg") + def suite(): testSuite = unittest.TestSuite() testSuite.addTest(TestAddition("test_add")) # testSuite.addTest(TestAddition("test_measurement")) + testSuite.addTest(TestAddition("test_query")) return testSuite diff --git a/silx/opencl/test/test_backprojection.py b/silx/opencl/test/test_backprojection.py index b2f2070..9dfdd3a 100644 --- a/silx/opencl/test/test_backprojection.py +++ b/silx/opencl/test/test_backprojection.py @@ -96,8 +96,9 @@ class TestFBP(unittest.TestCase): # Therefore, we cannot expect results to be the "same" (up to float32 # numerical error) self.tol = 5e-2 - if self.fbp.is_cpu: + if not(self.fbp._use_textures) or self.fbp.device.type == "CPU": # Precision is less when using CPU + # (either CPU textures or "manual" linear interpolation) self.tol *= 2 def tearDown(self): diff --git a/silx/opencl/test/test_convolution.py b/silx/opencl/test/test_convolution.py index 27cb8a9..7bceb0d 100644 --- a/silx/opencl/test/test_convolution.py +++ b/silx/opencl/test/test_convolution.py @@ -41,15 +41,18 @@ from itertools import product import numpy as np from silx.utils.testutils import parameterize from silx.image.utils import gaussian_kernel + try: from scipy.ndimage import convolve, convolve1d from scipy.misc import ascent + scipy_convolve = convolve scipy_convolve1d = convolve1d except ImportError: scipy_convolve = None import unittest -from ..common import ocl +from ..common import ocl, check_textures_availability + if ocl: import pyopencl as cl import pyopencl.array as parray @@ -59,7 +62,6 @@ logger = logging.getLogger(__name__) @unittest.skipUnless(ocl and scipy_convolve, "PyOpenCl/scipy is missing") class TestConvolution(unittest.TestCase): - @classmethod def setUpClass(cls): super(TestConvolution, cls).setUpClass() @@ -67,7 +69,7 @@ class TestConvolution(unittest.TestCase): cls.data1d = cls.image[0] cls.data2d = cls.image cls.data3d = np.tile(cls.image[224:-224, 224:-224], (62, 1, 1)) - cls.kernel1d = gaussian_kernel(1.) + cls.kernel1d = gaussian_kernel(1.0) cls.kernel2d = np.outer(cls.kernel1d, cls.kernel1d) cls.kernel3d = np.multiply.outer(cls.kernel2d, cls.kernel1d) cls.ctx = ocl.create_context() @@ -97,7 +99,7 @@ class TestConvolution(unittest.TestCase): ) return errmsg - def __init__(self, methodName='runTest', param=None): + def __init__(self, methodName="runTest", param=None): unittest.TestCase.__init__(self, methodName) self.param = param self.mode = param["boundary_handling"] @@ -107,32 +109,27 @@ class TestConvolution(unittest.TestCase): use_textures=%s, input_device=%s, output_device=%s """ % ( - self.mode, param["use_textures"], - param["input_on_device"], param["output_on_device"] + self.mode, + param["use_textures"], + param["input_on_device"], + param["output_on_device"], ) ) def instantiate_convol(self, shape, kernel, axes=None): - def is_fermi_device(dev): - try: - res = (dev.compute_capability_major_nv < 3) - except cl.LogicError: - res = False - except AttributeError: - res = False - return res - if (self.mode == "constant") and ( - not(self.param["use_textures"]) - or (self.ctx.devices[0].type == cl._cl.device_type.CPU) - or (is_fermi_device(self.ctx.devices[0])) + if self.mode == "constant": + if not (self.param["use_textures"]) or ( + self.param["use_textures"] + and not (check_textures_availability(self.ctx)) ): self.skipTest("mode=constant not implemented without textures") C = Convolution( - shape, kernel, + shape, + kernel, mode=self.mode, ctx=self.ctx, axes=axes, - extra_options={"dont_use_textures": not(self.param["use_textures"])} + extra_options={"dont_use_textures": not (self.param["use_textures"])}, ) return C @@ -142,13 +139,9 @@ class TestConvolution(unittest.TestCase): "test_separable_2D": (2, 1), "test_separable_3D": (3, 1), "test_nonseparable_2D": (2, 2), - "test_nonseparable_3D": (3, 3), - } - dim_data = { - 1: self.data1d, - 2: self.data2d, - 3: self.data3d + "test_nonseparable_3D": (3, 3), } + dim_data = {1: self.data1d, 2: self.data2d, 3: self.data3d} dim_kernel = { 1: self.kernel1d, 2: self.kernel2d, @@ -159,24 +152,26 @@ class TestConvolution(unittest.TestCase): def get_reference_function(self, test_name): ref_func = { - "test_1D": - lambda x, y : scipy_convolve1d(x, y, mode=self.mode), - "test_separable_2D": - lambda x, y : scipy_convolve1d( - scipy_convolve1d(x, y, mode=self.mode, axis=1), - y, mode=self.mode, axis=0 - ), - "test_separable_3D": - lambda x, y: scipy_convolve1d( - scipy_convolve1d( - scipy_convolve1d(x, y, mode=self.mode, axis=2), - y, mode=self.mode, axis=1), - y, mode=self.mode, axis=0 + "test_1D": lambda x, y: scipy_convolve1d(x, y, mode=self.mode), + "test_separable_2D": lambda x, y: scipy_convolve1d( + scipy_convolve1d(x, y, mode=self.mode, axis=1), + y, + mode=self.mode, + axis=0, + ), + "test_separable_3D": lambda x, y: scipy_convolve1d( + scipy_convolve1d( + scipy_convolve1d(x, y, mode=self.mode, axis=2), + y, + mode=self.mode, + axis=1, ), - "test_nonseparable_2D": - lambda x, y: scipy_convolve(x, y, mode=self.mode), - "test_nonseparable_3D": - lambda x, y : scipy_convolve(x, y, mode=self.mode), + y, + mode=self.mode, + axis=0, + ), + "test_nonseparable_2D": lambda x, y: scipy_convolve(x, y, mode=self.mode), + "test_nonseparable_3D": lambda x, y: scipy_convolve(x, y, mode=self.mode), } return ref_func[test_name] @@ -226,8 +221,8 @@ class TestConvolution(unittest.TestCase): data = self.data3d kernel = self.kernel2d conv = self.instantiate_convol(data.shape, kernel, axes=(0,)) - res = conv(data) # 3D - ref = scipy_convolve(data[0], kernel, mode=self.mode) # 2D + res = conv(data) # 3D + ref = scipy_convolve(data[0], kernel, mode=self.mode) # 2D std = np.std(res, axis=0) std_max = np.max(np.abs(std)) @@ -244,12 +239,9 @@ def test_convolution(): output_on_device_ = [True, False] testSuite = unittest.TestSuite() - param_vals = list(product( - boundary_handling_, - use_textures_, - input_on_device_, - output_on_device_ - )) + param_vals = list( + product(boundary_handling_, use_textures_, input_on_device_, output_on_device_) + ) for boundary_handling, use_textures, input_dev, output_dev in param_vals: testcase = parameterize( TestConvolution, @@ -258,17 +250,16 @@ def test_convolution(): "input_on_device": input_dev, "output_on_device": output_dev, "use_textures": use_textures, - } + }, ) testSuite.addTest(testcase) return testSuite - def suite(): testSuite = test_convolution() return testSuite -if __name__ == '__main__': +if __name__ == "__main__": unittest.main(defaultTest="suite") |