From 84bed8bbe77ed09671df0aad5f090d1ce4c66101 Mon Sep 17 00:00:00 2001 From: vorj <40021161+vorj@users.noreply.github.com> Date: Tue, 16 Apr 2019 15:46:22 +0900 Subject: [PATCH 1/8] implement clpy.backend.opencl.env.supports.cl_khr_fp16 --- clpy/backend/opencl/env.pyx | 92 +++++++++++++++++++++++++++++++++++++ 1 file changed, 92 insertions(+) diff --git a/clpy/backend/opencl/env.pyx b/clpy/backend/opencl/env.pyx index 01b8f8a50ac..a6e0ac1cb95 100644 --- a/clpy/backend/opencl/env.pyx +++ b/clpy/backend/opencl/env.pyx @@ -2,12 +2,18 @@ import atexit import locale import logging +import numpy import re +from clpy.backend.opencl import api from clpy.backend.opencl cimport api +from clpy.backend.opencl import exceptions from cython.view cimport array as cython_array from libc.stdlib cimport malloc +from libc.stdlib cimport free +from libc.string cimport memset +from libc.string cimport memcpy cdef interpret_versionstr(versionstr): version_detector = re.compile('''OpenCL (\d+)\.(\d+)''') @@ -135,6 +141,10 @@ for id in range(__num_devices): api.CreateCommandQueue(__contexts[id], __devices[id], 0) logging.info("SUCCESS") +cdef cl_char* __supports_cl_khr_fp16 = \ + malloc(sizeof(cl_char)*__num_devices) +memset(__supports_cl_khr_fp16, -1, sizeof(cl_char)*__num_devices) + ########################################## # Functions ########################################## @@ -162,6 +172,88 @@ cdef cl_device_id get_device(): global __current_device_id return __devices[__current_device_id] +cdef cl_char check_cl_khr_fp16(): + code = b''' + #pragma OPENCL EXTENSION cl_khr_fp16: enable + __kernel void check_cl_khr_fp16(__global half* v){ + ushort x = 0x5140; + *v = *(const half*)&x; + } + ''' + cdef size_t length + cdef char* src + cdef char* options + cdef cl_mem buf=NULL + cdef size_t global_work_size[3] + cdef size_t ptr + try: + device = __devices[__current_device_id] + context = __contexts[__current_device_id] + queue = __command_queues[__current_device_id] + + length = len(code) + src = malloc(sizeof(char)*length) + memcpy(src, code, length) + + program = api.CreateProgramWithSource( + context=context, + count=1, + strings=&src, + lengths=&length) + options = b'-cl-fp32-correctly-rounded-divide-sqrt' + api.BuildProgram( + program, + 1, + &device, + options, + NULL, + NULL) + kernel = api.CreateKernel( + program, b'check_cl_khr_fp16') + global_work_size[0] = 1 + global_work_size[1] = 0 + global_work_size[2] = 0 + v = numpy.array([0], dtype=numpy.float16) + ptr = v.ctypes.get_as_parameter().value + buf = api.CreateBuffer( + context, + CL_MEM_WRITE_ONLY, + v.nbytes, + NULL) + api.SetKernelArg( + kernel, + 0, + sizeof(cl_mem), + &buf) + api.EnqueueNDRangeKernel( + command_queue=queue, + kernel=kernel, + work_dim=1, + global_work_offset=NULL, + global_work_size=&global_work_size[0], + local_work_size=NULL) + api.EnqueueReadBuffer( + command_queue=queue, + buffer=buf, + blocking_read=api.BLOCKING, + offset=0, + cb=v.nbytes, + host_ptr=ptr) + result = 1 if v[0] == numpy.float16(42) else 0 + except exceptions.OpenCLRuntimeError: + result = 0 + if buf != NULL: + api.ReleaseMemObject(buf) + free(src) + return result + +cpdef cl_bool supports_cl_khr_fp16(): + global __current_device_id + if __supports_cl_khr_fp16[__current_device_id] == -1: + __supports_cl_khr_fp16[__current_device_id] = check_cl_khr_fp16() + return CL_TRUE if __supports_cl_khr_fp16[__current_device_id] == 1 \ + else CL_FALSE + def release(): """Release command_queue and context automatically.""" From 5caa98b05ab26dc548a65765913e502c329ba315 Mon Sep 17 00:00:00 2001 From: vorj <40021161+vorj@users.noreply.github.com> Date: Tue, 16 Apr 2019 15:50:42 +0900 Subject: [PATCH 2/8] define __CLPY_ENABLE_CL_KHR_FP16 when the device supports cl_khr_fp16 --- clpy/core/carray.pxi | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/clpy/core/carray.pxi b/clpy/core/carray.pxi index 20742a87a7c..3b9a2b26352 100644 --- a/clpy/core/carray.pxi +++ b/clpy/core/carray.pxi @@ -10,6 +10,7 @@ import warnings from clpy import backend from clpy.backend cimport function # from clpy.backend cimport runtime +import clpy.backend.opencl.api cimport clpy.backend.opencl.api cimport clpy.backend.opencl.utility import clpy.backend.opencl.env @@ -186,6 +187,9 @@ cpdef function.Module compile_with_cache( extra_source = _get_header_source() options += ('-I%s' % _get_header_dir_path(),) options += (' -cl-fp32-correctly-rounded-divide-sqrt', ) + if clpy.backend.opencl.env.supports_cl_khr_fp16() == \ + clpy.backend.opencl.api.TRUE: + options += (' -D__CLPY_ENABLE_CL_KHR_FP16', ) optionStr = functools.reduce(operator.add, options) device = clpy.backend.opencl.env.get_device() From 85fc8bb2e55adca9d2834bec4d0fa38c77254d71 Mon Sep 17 00:00:00 2001 From: vorj <40021161+vorj@users.noreply.github.com> Date: Tue, 16 Apr 2019 15:51:06 +0900 Subject: [PATCH 3/8] disable cl_khr_fp16 when __CLPY_ENABLE_CL_KHR_FP16 is not defined --- clpy/core/include/clpy/carray.clh | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/clpy/core/include/clpy/carray.clh b/clpy/core/include/clpy/carray.clh index 7efa290d0d4..0bb09cd1106 100644 --- a/clpy/core/include/clpy/carray.clh +++ b/clpy/core/include/clpy/carray.clh @@ -1,5 +1,7 @@ #pragma once +#ifdef __CLPY_ENABLE_CL_KHR_FP16 #pragma OPENCL EXTENSION cl_khr_fp16: enable +#endif // TODO: Implement common functions in OpenCL C #if 0 @@ -532,13 +534,16 @@ static void __clpy_end_print_out() __attribute__((annotate("clpy_end_print_out") #ifdef __ULTIMA __attribute__((annotate("clpy_no_mangle"))) static half convert_float_to_half(float x); #else +#ifdef __CLPY_ENABLE_CL_KHR_FP16 #include "fp16.clh" typedef half __clpy__half; #endif +#endif #ifdef __ULTIMA __attribute__((annotate("clpy_no_mangle"))) static half clpy_nextafter_fp16(half x1, half x2); #else +#ifdef __CLPY_ENABLE_CL_KHR_FP16 static int isnan_fp16(half x){ unsigned short const* x_raw = (unsigned short const*)&x; return (*x_raw & 0x7c00u) == 0x7c00u && (*x_raw & 0x03ffu) != 0x0000u; @@ -581,3 +586,4 @@ static half clpy_nextafter_fp16(half x1, half x2){ return *(half*)&ret_raw_; } #endif +#endif From 1af1deaa0fb4be0349bd5b4e498259689baede6e Mon Sep 17 00:00:00 2001 From: vorj <40021161+vorj@users.noreply.github.com> Date: Tue, 4 Feb 2020 00:43:28 +0900 Subject: [PATCH 4/8] add no_8bit_integer option to testing.helper.for_dtypes --- clpy/testing/__init__.py | 1 + clpy/testing/helper.py | 93 +++++++++++++++++++++++++++------------- 2 files changed, 65 insertions(+), 29 deletions(-) diff --git a/clpy/testing/__init__.py b/clpy/testing/__init__.py index 2e062322b1f..0c3cbb5ae76 100644 --- a/clpy/testing/__init__.py +++ b/clpy/testing/__init__.py @@ -17,6 +17,7 @@ from clpy.testing.bufio import readbuf # NOQA from clpy.testing.bufio import writebuf # NOQA from clpy.testing.helper import assert_warns # NOQA +from clpy.testing.helper import for_8bit_integer_dtypes # NOQA from clpy.testing.helper import for_all_dtypes # NOQA from clpy.testing.helper import for_all_dtypes_combination # NOQA from clpy.testing.helper import for_CF_orders # NOQA diff --git a/clpy/testing/helper.py b/clpy/testing/helper.py index 5680c812c35..57608377448 100644 --- a/clpy/testing/helper.py +++ b/clpy/testing/helper.py @@ -509,24 +509,31 @@ def test_func(self, *args, **kw): _complex_dtypes = (numpy.complex64, numpy.complex128) _regular_float_dtypes = (numpy.float64, numpy.float32) _float_dtypes = _regular_float_dtypes + (numpy.float16,) -_signed_dtypes = tuple(numpy.dtype(i).type for i in 'bhilq') -_unsigned_dtypes = tuple(numpy.dtype(i).type for i in 'BHILQ') +_signed_dtypes_without_int8 = tuple(numpy.dtype(i).type for i in 'hilq') +_signed_dtypes = (numpy.int8,) + _signed_dtypes_without_int8 +_unsigned_dtypes_without_uint8 = tuple(numpy.dtype(i).type for i in 'HILQ') +_unsigned_dtypes = (numpy.uint8,) + _unsigned_dtypes_without_uint8 +_int_dtypes_without_8bit = _signed_dtypes_without_int8 + \ + _unsigned_dtypes_without_uint8 _int_dtypes = _signed_dtypes + _unsigned_dtypes _int_bool_dtypes = _int_dtypes + (numpy.bool_,) _regular_dtypes = _regular_float_dtypes + _int_bool_dtypes -_dtypes = _float_dtypes + _int_bool_dtypes +_8bit_int_dtypes = tuple(numpy.dtype(i).type for i in 'bB') +_8bit_int_bool_dtypes = tuple(numpy.dtype(i).type for i in '?bB') -def _make_all_dtypes(no_float16, no_bool, no_complex): +def _make_all_dtypes(no_float16, no_bool, no_complex, no_8bit_integer): if no_float16: dtypes = _regular_float_dtypes else: dtypes = _float_dtypes - if no_bool: - dtypes += _int_dtypes - else: - dtypes += _int_bool_dtypes + dtypes += _int_dtypes_without_8bit + + if not no_8bit_integer: + dtypes += _8bit_int_dtypes + if not no_bool: + dtypes += (numpy.bool_,) # if not no_complex: # # TODO(LWisteria): Support complex in OpenCL @@ -536,7 +543,7 @@ def _make_all_dtypes(no_float16, no_bool, no_complex): def for_all_dtypes(name='dtype', no_float16=True, no_bool=False, - no_complex=False): + no_complex=False, no_8bit_integer=False): """Decorator that checks the fixture with all dtypes. Args: @@ -591,8 +598,10 @@ def for_all_dtypes(name='dtype', no_float16=True, no_bool=False, .. seealso:: :func:`clpy.testing.for_dtypes` """ - return for_dtypes(_make_all_dtypes(no_float16, no_bool, no_complex), - name=name) + return for_dtypes(_make_all_dtypes( + no_float16, no_bool, no_complex, + no_8bit_integer), + name=name) def for_float_dtypes(name='dtype', no_float16=True): @@ -615,7 +624,7 @@ def for_float_dtypes(name='dtype', no_float16=True): return for_dtypes(_float_dtypes, name=name) -def for_signed_dtypes(name='dtype'): +def for_signed_dtypes(name='dtype', no_int8=False): """Decorator that checks the fixture with signed dtypes. Args: @@ -627,10 +636,13 @@ def for_signed_dtypes(name='dtype'): .. seealso:: :func:`clpy.testing.for_dtypes`, :func:`clpy.testing.for_all_dtypes` """ - return for_dtypes(_signed_dtypes, name=name) + if no_int8: + return for_dtypes(_signed_dtypes_without_int8, name=name) + else: + return for_dtypes(_signed_dtypes, name=name) -def for_unsigned_dtypes(name='dtype'): +def for_unsigned_dtypes(name='dtype', no_uint8=False): """Decorator that checks the fixture with all dtypes. Args: @@ -643,10 +655,13 @@ def for_unsigned_dtypes(name='dtype'): .. seealso:: :func:`clpy.testing.for_dtypes`, :func:`clpy.testing.for_all_dtypes` """ - return for_dtypes(_unsigned_dtypes, name=name) + if no_uint8: + return for_dtypes(_unsigned_dtypes_without_uint8, name=name) + else: + return for_dtypes(_unsigned_dtypes, name=name) -def for_int_dtypes(name='dtype', no_bool=False): +def for_int_dtypes(name='dtype', no_bool=False, no_8bit_integer=False): """Decorator that checks the fixture with integer and optionally bool dtypes. Args: @@ -662,10 +677,17 @@ def for_int_dtypes(name='dtype', no_bool=False): .. seealso:: :func:`clpy.testing.for_dtypes`, :func:`clpy.testing.for_all_dtypes` """ - if no_bool: - return for_dtypes(_int_dtypes, name=name) + if no_8bit_integer: + return for_dtypes(_int_dtypes_without_8bit, name=name) else: - return for_dtypes(_int_bool_dtypes, name=name) + if no_bool: + return for_dtypes(_int_dtypes, name=name) + else: + return for_dtypes(_int_bool_dtypes, name=name) + + +def for_8bit_integer_dtypes(name='dtype'): + return for_dtypes(_8bit_int_bool_dtypes, name=name) def for_dtypes_combination(types, names=('dtype',), full=None): @@ -743,7 +765,7 @@ def test_func(self, *args, **kw): def for_all_dtypes_combination(names=('dtyes',), no_float16=True, no_bool=False, full=None, - no_complex=False): + no_complex=False, no_8bit_integer=False): """Decorator that checks the fixture with a product set of all dtypes. Args: @@ -761,11 +783,11 @@ def for_all_dtypes_combination(names=('dtyes',), .. seealso:: :func:`clpy.testing.for_dtypes_combination` """ - types = _make_all_dtypes(no_float16, no_bool, no_complex) + types = _make_all_dtypes(no_float16, no_bool, no_complex, no_8bit_integer) return for_dtypes_combination(types, names, full) -def for_signed_dtypes_combination(names=('dtype',), full=None): +def for_signed_dtypes_combination(names=('dtype',), no_int8=False, full=None): """Decorator for parameterized test w.r.t. the product set of signed dtypes. Args: @@ -777,10 +799,15 @@ def for_signed_dtypes_combination(names=('dtype',), full=None): .. seealso:: :func:`clpy.testing.for_dtypes_combination` """ - return for_dtypes_combination(_signed_dtypes, names=names, full=full) + if no_int8: + types = _signed_dtypes_without_int8 + else: + types = _signed_dtypes + return for_dtypes_combination(types, names=names, full=full) -def for_unsigned_dtypes_combination(names=('dtype',), full=None): +def for_unsigned_dtypes_combination(names=('dtype',), no_uint8=False, + full=None): """Decorator for parameterized test w.r.t. the product set of unsigned dtypes. Args: @@ -792,10 +819,15 @@ def for_unsigned_dtypes_combination(names=('dtype',), full=None): .. seealso:: :func:`clpy.testing.for_dtypes_combination` """ - return for_dtypes_combination(_unsigned_dtypes, names=names, full=full) + if no_uint8: + types = _unsigned_dtypes_without_uint8 + else: + types = _unsigned_dtypes + return for_dtypes_combination(types, names=names, full=full) -def for_int_dtypes_combination(names=('dtype',), no_bool=False, full=None): +def for_int_dtypes_combination(names=('dtype',), no_bool=False, + no_8bit_integer=False, full=None): """Decorator for parameterized test w.r.t. the product set of int and boolean. Args: @@ -809,10 +841,13 @@ def for_int_dtypes_combination(names=('dtype',), no_bool=False, full=None): .. seealso:: :func:`clpy.testing.for_dtypes_combination` """ - if no_bool: - types = _int_dtypes + if no_8bit_integer: + types = _int_dtypes_without_8bit else: - types = _int_bool_dtypes + if no_bool: + types = _int_dtypes + else: + types = _int_bool_dtypes return for_dtypes_combination(types, names, full) From 69c3a83bd46e77245930484a94904e210ad70f2f Mon Sep 17 00:00:00 2001 From: vorj <40021161+vorj@users.noreply.github.com> Date: Tue, 4 Feb 2020 00:45:12 +0900 Subject: [PATCH 5/8] add skip_when_disabled_cl_khr_fp16 --- clpy/testing/__init__.py | 1 + clpy/testing/attr.py | 8 ++++++++ 2 files changed, 9 insertions(+) diff --git a/clpy/testing/__init__.py b/clpy/testing/__init__.py index 0c3cbb5ae76..7718710521f 100644 --- a/clpy/testing/__init__.py +++ b/clpy/testing/__init__.py @@ -13,6 +13,7 @@ from clpy.testing.array import assert_array_max_ulp # NOQA from clpy.testing.attr import gpu # NOQA from clpy.testing.attr import multi_gpu # NOQA +from clpy.testing.attr import skip_when_disabled_cl_khr_fp16 # NOQA from clpy.testing.attr import slow # NOQA from clpy.testing.bufio import readbuf # NOQA from clpy.testing.bufio import writebuf # NOQA diff --git a/clpy/testing/attr.py b/clpy/testing/attr.py index c07daaa0b95..2e54cc1de34 100644 --- a/clpy/testing/attr.py +++ b/clpy/testing/attr.py @@ -1,3 +1,4 @@ +import clpy import os import unittest @@ -65,3 +66,10 @@ def gpu(f): check_available() return multi_gpu(1)(pytest.mark.gpu(f)) + + +def skip_when_disabled_cl_khr_fp16(f): + check_available() + return unittest.skipUnless( + clpy.backend.opencl.env.supports_cl_khr_fp16() == + clpy.backend.opencl.api.TRUE, "")(f) From 71522b0b0d8ebcf105eb803fbad004e2bc973e25 Mon Sep 17 00:00:00 2001 From: vorj <40021161+vorj@users.noreply.github.com> Date: Tue, 4 Feb 2020 01:15:58 +0900 Subject: [PATCH 6/8] separate 8bit integer test cases --- tests/clpy_tests/core_tests/test_fusion.py | 133 ++++++++++++++++-- .../core_tests/test_ndarray_complex_ops.py | 8 +- .../clpy_tests/math_tests/test_arithmetic.py | 30 +++- tests/clpy_tests/math_tests/test_explog.py | 26 +++- tests/clpy_tests/math_tests/test_floating.py | 11 +- .../clpy_tests/math_tests/test_hyperbolic.py | 12 +- tests/clpy_tests/math_tests/test_misc.py | 17 ++- tests/clpy_tests/math_tests/test_rounding.py | 22 ++- .../math_tests/test_trigonometric.py | 24 +++- 9 files changed, 262 insertions(+), 21 deletions(-) diff --git a/tests/clpy_tests/core_tests/test_fusion.py b/tests/clpy_tests/core_tests/test_fusion.py index 3e5f49b7ad8..0f7d25ec8fb 100644 --- a/tests/clpy_tests/core_tests/test_fusion.py +++ b/tests/clpy_tests/core_tests/test_fusion.py @@ -163,14 +163,14 @@ class TestFusionTrigonometric(unittest.TestCase): _multiprocess_can_split_ = True - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return (a,) - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() def check_binary(self, name, xp, dtype): @@ -178,6 +178,21 @@ def check_binary(self, name, xp, dtype): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return a, b + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + @fusion_default_array_equal() + def check_unary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return (a,) + + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + @fusion_default_array_equal() + def check_binary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + b = testing.shaped_reverse_arange((2, 3), xp, dtype) + return a, b + @testing.for_dtypes(['f', 'd']) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() @@ -187,12 +202,15 @@ def check_unary_unit(self, name, xp, dtype): def test_sin(self): self.check_unary('sin') + self.check_unary_8bit('sin') def test_cos(self): self.check_unary('cos') + self.check_unary_8bit('cos') def test_tan(self): self.check_unary('tan') + self.check_unary_8bit('tan') def test_arcsin(self): self.check_unary_unit('arcsin') @@ -202,18 +220,23 @@ def test_arccos(self): def test_arctan(self): self.check_unary('arctan') + self.check_unary_8bit('arctan') def test_arctan2(self): self.check_binary('arctan2') + self.check_binary_8bit('arctan2') def test_hypot(self): self.check_binary('hypot') + self.check_binary_8bit('hypot') def test_deg2rad(self): self.check_unary('deg2rad') + self.check_unary_8bit('deg2rad') def test_rad2deg(self): self.check_unary('rad2deg') + self.check_unary_8bit('rad2deg') @testing.gpu @@ -221,13 +244,20 @@ class TestFusionHyperbolic(unittest.TestCase): _multiprocess_can_split_ = True - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return (a,) + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + @fusion_default_array_equal() + def check_unary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return (a,) + @testing.for_dtypes(['f', 'd']) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() @@ -244,15 +274,19 @@ def check_unary_unit2(self, name, xp, dtype): def test_sinh(self): self.check_unary('sinh') + self.check_unary_8bit('sinh') def test_cosh(self): self.check_unary('cosh') + self.check_unary_8bit('cosh') def test_tanh(self): self.check_unary('tanh') + self.check_unary_8bit('tanh') def test_arcsinh(self): self.check_unary('arcsinh') + self.check_unary_8bit('arcsinh') def test_arccosh(self): self.check_unary_unit1('arccosh') @@ -266,37 +300,57 @@ class TestFusionRounding(unittest.TestCase): _multiprocess_can_split_ = True - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return (a,) - @testing.for_dtypes(['?', 'b', 'h', 'i', 'q', 'f', 'd']) + @testing.for_dtypes(['h', 'i', 'q', 'f', 'd']) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() def check_unary_negative(self, name, xp, dtype): a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) return (a,) + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + @fusion_default_array_equal() + def check_unary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return (a,) + + @testing.for_dtypes(['?', 'b']) + @testing.numpy_clpy_allclose(atol=1e-5) + @fusion_default_array_equal() + def check_unary_negative_8bit(self, name, xp, dtype): + a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) + return (a,) + def test_rint(self): self.check_unary('rint') + self.check_unary_8bit('rint') def test_rint_negative(self): self.check_unary_negative('rint') + self.check_unary_negative_8bit('rint') def test_floor(self): self.check_unary('floor') + self.check_unary_8bit('floor') def test_ceil(self): self.check_unary('ceil') + self.check_unary_8bit('ceil') def test_trunc(self): self.check_unary('trunc') + self.check_unary_8bit('trunc') def test_fix(self): self.check_unary('fix') + self.check_unary_8bit('fix') @testing.gpu @@ -304,14 +358,14 @@ class TestFusionExplog(unittest.TestCase): _multiprocess_can_split_ = True - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return (a,) - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() def check_binary(self, name, xp, dtype): @@ -319,35 +373,59 @@ def check_binary(self, name, xp, dtype): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return a, b + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + @fusion_default_array_equal() + def check_unary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return (a,) + + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + @fusion_default_array_equal() + def check_binary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + b = testing.shaped_reverse_arange((2, 3), xp, dtype) + return a, b + def test_exp(self): self.check_unary('exp') + self.check_unary_8bit('exp') def test_expm1(self): self.check_unary('expm1') + self.check_unary_8bit('expm1') def test_exp2(self): self.check_unary('exp2') + self.check_unary_8bit('exp2') def test_log(self): with testing.NumpyError(divide='ignore'): self.check_unary('log') + self.check_unary_8bit('log') def test_log10(self): with testing.NumpyError(divide='ignore'): self.check_unary('log10') + self.check_unary_8bit('log10') def test_log2(self): with testing.NumpyError(divide='ignore'): self.check_unary('log2') + self.check_unary_8bit('log2') def test_log1p(self): self.check_unary('log1p') + self.check_unary_8bit('log1p') def test_logaddexp(self): self.check_binary('logaddexp') + self.check_binary_8bit('logaddexp') def test_logaddexp2(self): self.check_binary('logaddexp2') + self.check_binary_8bit('logaddexp2') @testing.gpu @@ -362,7 +440,7 @@ def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return (a,) - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() def check_binary(self, name, xp, dtype): @@ -370,6 +448,14 @@ def check_binary(self, name, xp, dtype): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return a, b + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + @fusion_default_array_equal() + def check_binary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + b = testing.shaped_reverse_arange((2, 3), xp, dtype) + return a, b + @testing.for_float_dtypes(name='ftype') @testing.for_dtypes(['i', 'l'], name='itype') @testing.numpy_clpy_allclose() @@ -388,6 +474,7 @@ def test_signbit(self): def test_copysign(self): self.check_binary('copysign') + self.check_binary_8bit('copysign') @testing.for_float_dtypes() def test_frexp(self, dtype): @@ -408,6 +495,7 @@ def g(x): def test_nextafter(self): self.check_binary('nextafter') + self.check_binary_8bit('nextafter') @testing.gpu @@ -846,6 +934,32 @@ def g(x): return g(a) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_without_8bit(self, name, xp, dtype, no_bool=False): + if no_bool and numpy.dtype(dtype).char == '?': + return numpy.int_(0) + a = testing.shaped_arange((2, 3), xp, dtype) + + @clpy.fuse() + def g(x): + return getattr(clpy, name)(x) + + return g(a) + + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_8bit(self, name, xp, dtype, no_bool=False): + if no_bool and numpy.dtype(dtype).char == '?': + return numpy.int_(0) + a = testing.shaped_arange((2, 3), xp, dtype) + + @clpy.fuse() + def g(x): + return getattr(clpy, name)(x) + + return g(a) + @testing.for_all_dtypes(no_complex=True) @testing.numpy_clpy_allclose(atol=1e-5) def check_binary(self, name, xp, dtype, no_bool=False): @@ -897,7 +1011,8 @@ def g(x, y, z): @testing.with_requires('numpy>=1.11.2') def test_sqrt(self): # numpy.sqrt is broken in numpy<1.11.2 - self.check_unary('sqrt') + self.check_unary_without_8bit('sqrt') + self.check_unary_8bit('sqrt') def test_square(self): self.check_unary('square') diff --git a/tests/clpy_tests/core_tests/test_ndarray_complex_ops.py b/tests/clpy_tests/core_tests/test_ndarray_complex_ops.py index a92ba333f87..aa2359914e2 100644 --- a/tests/clpy_tests/core_tests/test_ndarray_complex_ops.py +++ b/tests/clpy_tests/core_tests/test_ndarray_complex_ops.py @@ -28,12 +28,18 @@ class TestAngle(unittest.TestCase): _multiprocess_can_split_ = True - @testing.for_all_dtypes() + @testing.for_all_dtypes(no_8bit_integer=True) @testing.numpy_clpy_array_almost_equal() def test_angle(self, xp, dtype): x = testing.shaped_arange((2, 3), xp, dtype) return xp.angle(x) + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_array_almost_equal() + def test_angle_8bit(self, xp, dtype): + x = testing.shaped_arange((2, 3), xp, dtype) + return xp.angle(x) + @testing.gpu class TestRealImag(unittest.TestCase): diff --git a/tests/clpy_tests/math_tests/test_arithmetic.py b/tests/clpy_tests/math_tests/test_arithmetic.py index a22348cf6a7..7e0c8d99648 100644 --- a/tests/clpy_tests/math_tests/test_arithmetic.py +++ b/tests/clpy_tests/math_tests/test_arithmetic.py @@ -17,6 +17,18 @@ def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) + @testing.for_all_dtypes(no_8bit_integer=True) + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_without_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return getattr(xp, name)(a) + + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return getattr(xp, name)(a) + @testing.for_all_dtypes() @testing.numpy_clpy_allclose(atol=1e-4) def check_binary(self, name, xp, dtype, no_complex=False, no_bool=False): @@ -34,6 +46,18 @@ def check_unary_negative(self, name, xp, dtype): a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) return getattr(xp, name)(a) + @testing.for_dtypes(['h', 'i', 'q', 'f', 'd']) + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_negative_without_8bit(self, name, xp, dtype): + a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) + return getattr(xp, name)(a) + + @testing.for_dtypes(['?', 'b']) + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_negative_8bit(self, name, xp, dtype): + a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) + return getattr(xp, name)(a) + @testing.for_dtypes(['?', 'b', 'h', 'i', 'q', 'f', 'd']) @testing.numpy_clpy_allclose(atol=1e-4) def check_binary_negative(self, name, xp, dtype): @@ -151,9 +175,11 @@ def test_conj(self): self.check_raises_with_numpy_input(1, 'conj') def test_angle(self): - self.check_unary('angle') - self.check_unary_negative('angle') + self.check_unary_without_8bit('angle') + self.check_unary_negative_without_8bit('angle') self.check_raises_with_numpy_input(1, 'angle') + self.check_unary_8bit('angle') + self.check_unary_negative_8bit('angle') def test_real(self): self.check_unary('real') diff --git a/tests/clpy_tests/math_tests/test_explog.py b/tests/clpy_tests/math_tests/test_explog.py index c83aa0c338a..8f0ecb7364c 100644 --- a/tests/clpy_tests/math_tests/test_explog.py +++ b/tests/clpy_tests/math_tests/test_explog.py @@ -10,7 +10,7 @@ class TestExplog(unittest.TestCase): _multiprocess_can_split_ = True - @testing.for_all_dtypes() + @testing.for_all_dtypes(no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) def check_unary(self, name, xp, dtype, no_complex=False): if no_complex: @@ -19,7 +19,7 @@ def check_unary(self, name, xp, dtype, no_complex=False): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) - @testing.for_all_dtypes() + @testing.for_all_dtypes(no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) def check_binary(self, name, xp, dtype, no_complex=False): if no_complex: @@ -29,32 +29,54 @@ def check_binary(self, name, xp, dtype, no_complex=False): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return getattr(xp, name)(a, b) + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return getattr(xp, name)(a) + + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + def check_binary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + b = testing.shaped_reverse_arange((2, 3), xp, dtype) + return getattr(xp, name)(a, b) + def test_exp(self): self.check_unary('exp') + self.check_unary_8bit('exp') def test_expm1(self): self.check_unary('expm1', no_complex=True) + self.check_unary_8bit('expm1') def test_exp2(self): self.check_unary('exp2') + self.check_unary_8bit('exp2') def test_log(self): with testing.NumpyError(divide='ignore'): self.check_unary('log') + self.check_unary_8bit('log') def test_log10(self): with testing.NumpyError(divide='ignore'): self.check_unary('log10') + self.check_unary_8bit('log10') def test_log2(self): with testing.NumpyError(divide='ignore'): self.check_unary('log2', no_complex=True) + self.check_unary_8bit('log2') def test_log1p(self): self.check_unary('log1p', no_complex=True) + self.check_unary_8bit('log1p') def test_logaddexp(self): self.check_binary('logaddexp', no_complex=True) + self.check_binary_8bit('logaddexp') def test_logaddexp2(self): self.check_binary('logaddexp2', no_complex=True) + self.check_binary_8bit('logaddexp2') diff --git a/tests/clpy_tests/math_tests/test_floating.py b/tests/clpy_tests/math_tests/test_floating.py index ed64b89e480..f07c010b66c 100644 --- a/tests/clpy_tests/math_tests/test_floating.py +++ b/tests/clpy_tests/math_tests/test_floating.py @@ -19,18 +19,26 @@ def check_unary(self, name, xp, dtype, no_complex=False): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) def check_binary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) b = testing.shaped_reverse_arange((2, 3), xp, dtype) return getattr(xp, name)(a, b) + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + def check_binary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + b = testing.shaped_reverse_arange((2, 3), xp, dtype) + return getattr(xp, name)(a, b) + def test_signbit(self): self.check_unary('signbit', no_complex=True) def test_copysign(self): self.check_binary('copysign') + self.check_binary_8bit('copysign') @testing.for_float_dtypes(name='ftype') @testing.for_dtypes(['i', 'l'], name='itype') @@ -54,3 +62,4 @@ def test_frexp(self, dtype): def test_nextafter(self): self.check_binary('nextafter') + self.check_binary_8bit('nextafter') diff --git a/tests/clpy_tests/math_tests/test_hyperbolic.py b/tests/clpy_tests/math_tests/test_hyperbolic.py index a77a17ae982..33f17cdf070 100644 --- a/tests/clpy_tests/math_tests/test_hyperbolic.py +++ b/tests/clpy_tests/math_tests/test_hyperbolic.py @@ -8,12 +8,18 @@ class TestHyperbolic(unittest.TestCase): _multiprocess_can_split_ = True - @testing.for_all_dtypes() + @testing.for_all_dtypes(no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return getattr(xp, name)(a) + @testing.for_dtypes(['f', 'd']) @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_unit(self, name, xp, dtype): @@ -22,15 +28,19 @@ def check_unary_unit(self, name, xp, dtype): def test_sinh(self): self.check_unary('sinh') + self.check_unary_8bit('sinh') def test_cosh(self): self.check_unary('cosh') + self.check_unary_8bit('cosh') def test_tanh(self): self.check_unary('tanh') + self.check_unary_8bit('tanh') def test_arcsinh(self): self.check_unary('arcsinh') + self.check_unary_8bit('arcsinh') @testing.for_dtypes(['f', 'd']) @testing.numpy_clpy_allclose(atol=1e-5) diff --git a/tests/clpy_tests/math_tests/test_misc.py b/tests/clpy_tests/math_tests/test_misc.py index 16183f91b43..84511fad268 100644 --- a/tests/clpy_tests/math_tests/test_misc.py +++ b/tests/clpy_tests/math_tests/test_misc.py @@ -20,6 +20,20 @@ def check_unary(self, name, xp, dtype, no_bool=False): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_without_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return getattr(xp, name)(a) + + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_8bit(self, name, xp, dtype): + if numpy.dtype(dtype).char == '?': + return numpy.int_(0) + a = testing.shaped_arange((2, 3), xp, dtype) + return getattr(xp, name)(a) + @testing.for_all_dtypes(no_complex=True) @testing.numpy_clpy_allclose(atol=1e-5) def check_binary(self, name, xp, dtype, no_bool=False): @@ -85,7 +99,8 @@ def test_clip2(self, xp, dtype): @testing.with_requires('numpy>=1.11.2') def test_sqrt(self): # numpy.sqrt is broken in numpy<1.11.2 - self.check_unary('sqrt') + self.check_unary_without_8bit('sqrt') + self.check_unary_8bit('sqrt') def test_square(self): self.check_unary('square') diff --git a/tests/clpy_tests/math_tests/test_rounding.py b/tests/clpy_tests/math_tests/test_rounding.py index 4626e22b039..14e6b137531 100644 --- a/tests/clpy_tests/math_tests/test_rounding.py +++ b/tests/clpy_tests/math_tests/test_rounding.py @@ -8,32 +8,50 @@ class TestRounding(unittest.TestCase): _multiprocess_can_split_ = True - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) - @testing.for_dtypes(['?', 'b', 'h', 'i', 'q', 'f', 'd']) + @testing.for_dtypes(['h', 'i', 'q', 'f', 'd']) @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_negative(self, name, xp, dtype): a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) return getattr(xp, name)(a) + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return getattr(xp, name)(a) + + @testing.for_dtypes(['?', 'b']) + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_negative_8bit(self, name, xp, dtype): + a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) + return getattr(xp, name)(a) + def test_rint(self): self.check_unary('rint') + self.check_unary_8bit('rint') def test_rint_negative(self): self.check_unary_negative('rint') + self.check_unary_negative_8bit('rint') def test_floor(self): self.check_unary('floor') + self.check_unary_8bit('floor') def test_ceil(self): self.check_unary('ceil') + self.check_unary_8bit('ceil') def test_trunc(self): self.check_unary('trunc') + self.check_unary_8bit('trunc') def test_fix(self): self.check_unary('fix') + self.check_unary_8bit('fix') diff --git a/tests/clpy_tests/math_tests/test_trigonometric.py b/tests/clpy_tests/math_tests/test_trigonometric.py index edbed20de31..b646c253665 100644 --- a/tests/clpy_tests/math_tests/test_trigonometric.py +++ b/tests/clpy_tests/math_tests/test_trigonometric.py @@ -8,19 +8,32 @@ class TestTrigonometric(unittest.TestCase): _multiprocess_can_split_ = True - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) - @testing.for_all_dtypes(no_complex=True) + @testing.for_all_dtypes(no_complex=True, no_8bit_integer=True) @testing.numpy_clpy_allclose(atol=1e-5) def check_binary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) b = testing.shaped_reverse_arange((2, 3), xp, dtype) return getattr(xp, name)(a, b) + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + def check_unary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + return getattr(xp, name)(a) + + @testing.for_8bit_integer_dtypes() + @testing.numpy_clpy_allclose(atol=1e-5) + def check_binary_8bit(self, name, xp, dtype): + a = testing.shaped_arange((2, 3), xp, dtype) + b = testing.shaped_reverse_arange((2, 3), xp, dtype) + return getattr(xp, name)(a, b) + @testing.for_dtypes(['f', 'd']) @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_unit(self, name, xp, dtype): @@ -29,12 +42,15 @@ def check_unary_unit(self, name, xp, dtype): def test_sin(self): self.check_unary('sin') + self.check_unary_8bit('sin') def test_cos(self): self.check_unary('cos') + self.check_unary_8bit('cos') def test_tan(self): self.check_unary('tan') + self.check_unary_8bit('tan') def test_arcsin(self): self.check_unary_unit('arcsin') @@ -44,15 +60,19 @@ def test_arccos(self): def test_arctan(self): self.check_unary('arctan') + self.check_unary_8bit('arctan') def test_arctan2(self): self.check_binary('arctan2') + self.check_binary_8bit('arctan2') def test_hypot(self): self.check_binary('hypot') + self.check_binary_8bit('hypot') def test_deg2rad(self): self.check_unary('deg2rad') + self.check_unary_8bit('deg2rad') def test_rad2deg(self): self.check_unary('rad2deg') From 390cbb0665133da44ed2b74fba05bcd7cc55a13d Mon Sep 17 00:00:00 2001 From: vorj <40021161+vorj@users.noreply.github.com> Date: Tue, 4 Feb 2020 01:16:51 +0900 Subject: [PATCH 7/8] add @testing.skip_when_disabled_cl_khr_fp16 to 8bit integer test cases --- tests/clpy_tests/core_tests/test_fusion.py | 9 +++++++++ tests/clpy_tests/core_tests/test_ndarray_complex_ops.py | 1 + tests/clpy_tests/math_tests/test_arithmetic.py | 2 ++ tests/clpy_tests/math_tests/test_explog.py | 2 ++ tests/clpy_tests/math_tests/test_floating.py | 1 + tests/clpy_tests/math_tests/test_hyperbolic.py | 1 + tests/clpy_tests/math_tests/test_misc.py | 1 + tests/clpy_tests/math_tests/test_rounding.py | 2 ++ tests/clpy_tests/math_tests/test_trigonometric.py | 2 ++ 9 files changed, 21 insertions(+) diff --git a/tests/clpy_tests/core_tests/test_fusion.py b/tests/clpy_tests/core_tests/test_fusion.py index 0f7d25ec8fb..d65fd4c51ba 100644 --- a/tests/clpy_tests/core_tests/test_fusion.py +++ b/tests/clpy_tests/core_tests/test_fusion.py @@ -178,6 +178,7 @@ def check_binary(self, name, xp, dtype): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return a, b + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() @@ -185,6 +186,7 @@ def check_unary_8bit(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return (a,) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() @@ -251,6 +253,7 @@ def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return (a,) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() @@ -314,6 +317,7 @@ def check_unary_negative(self, name, xp, dtype): a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) return (a,) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() @@ -321,6 +325,7 @@ def check_unary_8bit(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return (a,) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_dtypes(['?', 'b']) @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() @@ -373,6 +378,7 @@ def check_binary(self, name, xp, dtype): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return a, b + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() @@ -380,6 +386,7 @@ def check_unary_8bit(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return (a,) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() @@ -448,6 +455,7 @@ def check_binary(self, name, xp, dtype): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return a, b + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) @fusion_default_array_equal() @@ -947,6 +955,7 @@ def g(x): return g(a) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_8bit(self, name, xp, dtype, no_bool=False): diff --git a/tests/clpy_tests/core_tests/test_ndarray_complex_ops.py b/tests/clpy_tests/core_tests/test_ndarray_complex_ops.py index aa2359914e2..88a78bc22af 100644 --- a/tests/clpy_tests/core_tests/test_ndarray_complex_ops.py +++ b/tests/clpy_tests/core_tests/test_ndarray_complex_ops.py @@ -34,6 +34,7 @@ def test_angle(self, xp, dtype): x = testing.shaped_arange((2, 3), xp, dtype) return xp.angle(x) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_array_almost_equal() def test_angle_8bit(self, xp, dtype): diff --git a/tests/clpy_tests/math_tests/test_arithmetic.py b/tests/clpy_tests/math_tests/test_arithmetic.py index 7e0c8d99648..f1e956cacb3 100644 --- a/tests/clpy_tests/math_tests/test_arithmetic.py +++ b/tests/clpy_tests/math_tests/test_arithmetic.py @@ -23,6 +23,7 @@ def check_unary_without_8bit(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_8bit(self, name, xp, dtype): @@ -52,6 +53,7 @@ def check_unary_negative_without_8bit(self, name, xp, dtype): a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) return getattr(xp, name)(a) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_dtypes(['?', 'b']) @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_negative_8bit(self, name, xp, dtype): diff --git a/tests/clpy_tests/math_tests/test_explog.py b/tests/clpy_tests/math_tests/test_explog.py index 8f0ecb7364c..43f9e9cc3a0 100644 --- a/tests/clpy_tests/math_tests/test_explog.py +++ b/tests/clpy_tests/math_tests/test_explog.py @@ -29,12 +29,14 @@ def check_binary(self, name, xp, dtype, no_complex=False): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return getattr(xp, name)(a, b) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_8bit(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) def check_binary_8bit(self, name, xp, dtype): diff --git a/tests/clpy_tests/math_tests/test_floating.py b/tests/clpy_tests/math_tests/test_floating.py index f07c010b66c..4f9d5eb905e 100644 --- a/tests/clpy_tests/math_tests/test_floating.py +++ b/tests/clpy_tests/math_tests/test_floating.py @@ -26,6 +26,7 @@ def check_binary(self, name, xp, dtype): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return getattr(xp, name)(a, b) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) def check_binary_8bit(self, name, xp, dtype): diff --git a/tests/clpy_tests/math_tests/test_hyperbolic.py b/tests/clpy_tests/math_tests/test_hyperbolic.py index 33f17cdf070..1216404884a 100644 --- a/tests/clpy_tests/math_tests/test_hyperbolic.py +++ b/tests/clpy_tests/math_tests/test_hyperbolic.py @@ -14,6 +14,7 @@ def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_8bit(self, name, xp, dtype): diff --git a/tests/clpy_tests/math_tests/test_misc.py b/tests/clpy_tests/math_tests/test_misc.py index 84511fad268..fc3a2881065 100644 --- a/tests/clpy_tests/math_tests/test_misc.py +++ b/tests/clpy_tests/math_tests/test_misc.py @@ -26,6 +26,7 @@ def check_unary_without_8bit(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_8bit(self, name, xp, dtype): diff --git a/tests/clpy_tests/math_tests/test_rounding.py b/tests/clpy_tests/math_tests/test_rounding.py index 14e6b137531..cb3f1ca73ba 100644 --- a/tests/clpy_tests/math_tests/test_rounding.py +++ b/tests/clpy_tests/math_tests/test_rounding.py @@ -20,12 +20,14 @@ def check_unary_negative(self, name, xp, dtype): a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) return getattr(xp, name)(a) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_8bit(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_dtypes(['?', 'b']) @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_negative_8bit(self, name, xp, dtype): diff --git a/tests/clpy_tests/math_tests/test_trigonometric.py b/tests/clpy_tests/math_tests/test_trigonometric.py index b646c253665..79d3efa6629 100644 --- a/tests/clpy_tests/math_tests/test_trigonometric.py +++ b/tests/clpy_tests/math_tests/test_trigonometric.py @@ -21,12 +21,14 @@ def check_binary(self, name, xp, dtype): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return getattr(xp, name)(a, b) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) def check_unary_8bit(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) + @testing.skip_when_disabled_cl_khr_fp16 @testing.for_8bit_integer_dtypes() @testing.numpy_clpy_allclose(atol=1e-5) def check_binary_8bit(self, name, xp, dtype): From 43e8d6fdff36cd34014fe088a3bbdadaec7e1416 Mon Sep 17 00:00:00 2001 From: Yoshiki Imaizumi Date: Tue, 4 Feb 2020 14:21:47 +0900 Subject: [PATCH 8/8] fix warning error of importing imp --- tests/clpy_tests/testing_tests/test_helper.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/tests/clpy_tests/testing_tests/test_helper.py b/tests/clpy_tests/testing_tests/test_helper.py index 1f7e0b6ff66..b900ce0f056 100644 --- a/tests/clpy_tests/testing_tests/test_helper.py +++ b/tests/clpy_tests/testing_tests/test_helper.py @@ -231,6 +231,10 @@ def correct_failure(self, xp, dtype1, dtype2): return xp.array(-2, dtype=numpy.float32) def test_correct_failure(self): + # TODO(vorj): remove these warning suppressings + # when we will support numpy >= 1.15.0 + import warnings + warnings.filterwarnings("ignore", category=DeprecationWarning) numpy.testing.assert_raises_regex( AssertionError, 'mismatch 100.0%', self.correct_failure)