From 2dc6f22d9ecc048dd31a2c88110f2177cc8d6db4 Mon Sep 17 00:00:00 2001 From: SolidAhmad Date: Sat, 3 Dec 2022 08:47:16 -0700 Subject: [PATCH 1/9] fix skip test when no GPU available --- tests/test_core.py | 31 +++++++++++++------------------ 1 file changed, 13 insertions(+), 18 deletions(-) diff --git a/tests/test_core.py b/tests/test_core.py index a8e7cba8e..15a38a4ed 100644 --- a/tests/test_core.py +++ b/tests/test_core.py @@ -13,6 +13,17 @@ if cuda.is_available(): from stumpy.core import _gpu_searchsorted_left, _gpu_searchsorted_right + + @cuda.jit("(f8[:, :], f8[:], i8[:], i8, b1, i8[:])") + def _gpu_searchsorted_kernel(a, v, bfs, nlevel, is_left, idx): + # A wrapper kernel for calling device function _gpu_searchsorted_left/right. + i = cuda.grid(1) + if i < a.shape[0]: + if is_left: + idx[i] = _gpu_searchsorted_left(a[i], v[i], bfs, nlevel) + else: + idx[i] = _gpu_searchsorted_right(a[i], v[i], bfs, nlevel) + else: # pragma: no cover from stumpy.core import ( _gpu_searchsorted_left_driver_not_found as _gpu_searchsorted_left, @@ -26,22 +37,6 @@ except ModuleNotFoundError: from numba.core.errors import NumbaPerformanceWarning -TEST_THREADS_PER_BLOCK = 10 - -if not cuda.is_available(): # pragma: no cover - pytest.skip("Skipping Tests No GPUs Available", allow_module_level=True) - - -@cuda.jit("(f8[:, :], f8[:], i8[:], i8, b1, i8[:])") -def _gpu_searchsorted_kernel(a, v, bfs, nlevel, is_left, idx): - # A wrapper kernel for calling device function _gpu_searchsorted_left/right. - i = cuda.grid(1) - if i < a.shape[0]: - if is_left: - idx[i] = _gpu_searchsorted_left(a[i], v[i], bfs, nlevel) - else: - idx[i] = _gpu_searchsorted_right(a[i], v[i], bfs, nlevel) - def naive_rolling_window_dot_product(Q, T): window = len(Q) @@ -1400,10 +1395,10 @@ def test_find_matches_maxmatch(): @pytest.mark.filterwarnings("ignore", category=NumbaPerformanceWarning) -@patch("stumpy.config.STUMPY_THREADS_PER_BLOCK", TEST_THREADS_PER_BLOCK) +@patch("stumpy.config.STUMPY_THREADS_PER_BLOCK", 10) def test_gpu_searchsorted(): if not cuda.is_available(): # pragma: no cover - pytest.skip("Skipping Tests No GPUs Available", allow_module_level=True) + pytest.skip("Skipping Tests No GPUs Available") n = 3 * config.STUMPY_THREADS_PER_BLOCK + 1 V = np.empty(n, dtype=np.float64) From 36281d4d22359c5ef8010967156b36a3e88a8bb7 Mon Sep 17 00:00:00 2001 From: SolidAhmad Date: Sun, 4 Dec 2022 08:50:12 -0700 Subject: [PATCH 2/9] add conditional import in __init__ and update files --- stumpy/__init__.py | 39 +++++++++++++++++++++++++++++++++++++++ tests/test_core.py | 9 +-------- 2 files changed, 40 insertions(+), 8 deletions(-) diff --git a/stumpy/__init__.py b/stumpy/__init__.py index 6ce63705e..9f7693fd5 100644 --- a/stumpy/__init__.py +++ b/stumpy/__init__.py @@ -38,6 +38,8 @@ from .gpu_aampdist import gpu_aampdist # noqa: F401 from .gpu_stimp import gpu_stimp # noqa: F401 from .gpu_aamp_stimp import gpu_aamp_stimp # noqa: F401 + from .core import _gpu_searchsorted_left # noqa: F401 + from .core import _gpu_searchsorted_right # noqa: F401 else: # pragma: no cover from .core import _gpu_stump_driver_not_found as gpu_stump # noqa: F401 from .core import _gpu_aamp_driver_not_found as gpu_aamp # noqa: F401 @@ -49,6 +51,13 @@ from .core import _gpu_aampdist_driver_not_found as gpu_aampdist # noqa: F401 from .core import _gpu_stimp_driver_not_found as gpu_stimp # noqa: F401 from .core import _gpu_aamp_stimp_driver_not_found as gpu_aamp_stimp # noqa: F401 + from .core import ( + _gpu_searchsorted_left_driver_not_found as _gpu_searchsorted_left, + ) # noqa: F401 + from .core import ( + _gpu_searchsorted_right_driver_not_found as _gpu_searchsorted_right, + ) # noqa: F401 + import ast import pathlib @@ -174,6 +183,36 @@ if cd.name == "gpu_aamp_stimp": gpu_aamp_stimp.__doc__ = ast.get_docstring(cd) + # Fix _gpu_searchsorted_left Docs + _gpu_searchsorted_left.__doc__ = "" + filepath = pathlib.Path(__file__).parent / "core.py" + + file_contents = "" + with open(filepath, encoding="utf8") as f: + file_contents = f.read() + module = ast.parse(file_contents) + function_definitions = [ + node for node in module.body if isinstance(node, ast.FunctionDef) + ] + for fd in function_definitions: + if fd.name == "_gpu_searchsorted_left": + _gpu_searchsorted_left.__doc__ = ast.get_docstring(fd) + + # Fix _gpu_searchsorted_right Docs + _gpu_searchsorted_right.__doc__ = "" + filepath = pathlib.Path(__file__).parent / "core.py" + + file_contents = "" + with open(filepath, encoding="utf8") as f: + file_contents = f.read() + module = ast.parse(file_contents) + function_definitions = [ + node for node in module.body if isinstance(node, ast.FunctionDef) + ] + for fd in function_definitions: + if fd.name == "_gpu_searchsorted_right": + _gpu_searchsorted_right.__doc__ = ast.get_docstring(fd) + try: _dist = get_distribution("stumpy") # Normalize case for Windows systems diff --git a/tests/test_core.py b/tests/test_core.py index 15a38a4ed..a1a2ddd78 100644 --- a/tests/test_core.py +++ b/tests/test_core.py @@ -4,6 +4,7 @@ import pandas as pd from scipy.spatial.distance import cdist from stumpy import core, config +from stumpy.core import _gpu_searchsorted_left, _gpu_searchsorted_right import pytest from unittest.mock import patch import os @@ -12,7 +13,6 @@ import naive if cuda.is_available(): - from stumpy.core import _gpu_searchsorted_left, _gpu_searchsorted_right @cuda.jit("(f8[:, :], f8[:], i8[:], i8, b1, i8[:])") def _gpu_searchsorted_kernel(a, v, bfs, nlevel, is_left, idx): @@ -24,13 +24,6 @@ def _gpu_searchsorted_kernel(a, v, bfs, nlevel, is_left, idx): else: idx[i] = _gpu_searchsorted_right(a[i], v[i], bfs, nlevel) -else: # pragma: no cover - from stumpy.core import ( - _gpu_searchsorted_left_driver_not_found as _gpu_searchsorted_left, - ) - from stumpy.core import ( - _gpu_searchsorted_right_driver_not_found as _gpu_searchsorted_right, - ) try: from numba.errors import NumbaPerformanceWarning From 4fd9444460f601fa3f84d5a57e3ed0fa5cadea44 Mon Sep 17 00:00:00 2001 From: SolidAhmad Date: Sun, 4 Dec 2022 09:06:06 -0700 Subject: [PATCH 3/9] change import style for gpu device functions --- stumpy/gpu_aamp.py | 3 ++- stumpy/gpu_stump.py | 3 ++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/stumpy/gpu_aamp.py b/stumpy/gpu_aamp.py index bf431f71e..9d2fa1caf 100644 --- a/stumpy/gpu_aamp.py +++ b/stumpy/gpu_aamp.py @@ -10,6 +10,7 @@ from numba import cuda from . import core, config +from .core import _gpu_searchsorted_right logger = logging.getLogger(__name__) @@ -186,7 +187,7 @@ def _compute_and_update_PI_kernel( indices_R[j] = i if p_norm < profile[j, -1]: - idx = core._gpu_searchsorted_right(profile[j], p_norm, bfs, nlevel) + idx = _gpu_searchsorted_right(profile[j], p_norm, bfs, nlevel) for g in range(k - 1, idx, -1): profile[j, g] = profile[j, g - 1] indices[j, g] = indices[j, g - 1] diff --git a/stumpy/gpu_stump.py b/stumpy/gpu_stump.py index 5538e8696..42f054913 100644 --- a/stumpy/gpu_stump.py +++ b/stumpy/gpu_stump.py @@ -10,6 +10,7 @@ from numba import cuda from . import core, config +from .core import _gpu_searchsorted_right from .gpu_aamp import gpu_aamp logger = logging.getLogger(__name__) @@ -194,7 +195,7 @@ def _compute_and_update_PI_kernel( indices_R[j] = i if p_norm < profile[j, -1]: - idx = core._gpu_searchsorted_right(profile[j], p_norm, bfs, nlevel) + idx = _gpu_searchsorted_right(profile[j], p_norm, bfs, nlevel) for g in range(k - 1, idx, -1): profile[j, g] = profile[j, g - 1] indices[j, g] = indices[j, g - 1] From a06aa474f6081a7d648cd958e3f4d7b104b13896 Mon Sep 17 00:00:00 2001 From: SolidAhmad Date: Sun, 4 Dec 2022 13:08:14 -0700 Subject: [PATCH 4/9] Handle importing gpu device func when there is no GPU --- stumpy/__init__.py | 11 +++-------- 1 file changed, 3 insertions(+), 8 deletions(-) diff --git a/stumpy/__init__.py b/stumpy/__init__.py index 9f7693fd5..766df1e4d 100644 --- a/stumpy/__init__.py +++ b/stumpy/__init__.py @@ -38,9 +38,8 @@ from .gpu_aampdist import gpu_aampdist # noqa: F401 from .gpu_stimp import gpu_stimp # noqa: F401 from .gpu_aamp_stimp import gpu_aamp_stimp # noqa: F401 - from .core import _gpu_searchsorted_left # noqa: F401 - from .core import _gpu_searchsorted_right # noqa: F401 else: # pragma: no cover + from . import core from .core import _gpu_stump_driver_not_found as gpu_stump # noqa: F401 from .core import _gpu_aamp_driver_not_found as gpu_aamp # noqa: F401 from .core import _gpu_ostinato_driver_not_found as gpu_ostinato # noqa: F401 @@ -51,12 +50,8 @@ from .core import _gpu_aampdist_driver_not_found as gpu_aampdist # noqa: F401 from .core import _gpu_stimp_driver_not_found as gpu_stimp # noqa: F401 from .core import _gpu_aamp_stimp_driver_not_found as gpu_aamp_stimp # noqa: F401 - from .core import ( - _gpu_searchsorted_left_driver_not_found as _gpu_searchsorted_left, - ) # noqa: F401 - from .core import ( - _gpu_searchsorted_right_driver_not_found as _gpu_searchsorted_right, - ) # noqa: F401 + core._gpu_searchsorted_left = core._gpu_searchsorted_left_driver_not_found + core._gpu_searchsorted_right = core._gpu_searchsorted_right_driver_not_found import ast import pathlib From d2621b801db61bfccf4997139eada161b4a7112c Mon Sep 17 00:00:00 2001 From: SolidAhmad Date: Sun, 4 Dec 2022 13:17:41 -0700 Subject: [PATCH 5/9] replace gpu dummy docstring with actual one when no GPU --- stumpy/__init__.py | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/stumpy/__init__.py b/stumpy/__init__.py index 766df1e4d..56e8974d2 100644 --- a/stumpy/__init__.py +++ b/stumpy/__init__.py @@ -50,6 +50,7 @@ from .core import _gpu_aampdist_driver_not_found as gpu_aampdist # noqa: F401 from .core import _gpu_stimp_driver_not_found as gpu_stimp # noqa: F401 from .core import _gpu_aamp_stimp_driver_not_found as gpu_aamp_stimp # noqa: F401 + core._gpu_searchsorted_left = core._gpu_searchsorted_left_driver_not_found core._gpu_searchsorted_right = core._gpu_searchsorted_right_driver_not_found @@ -179,7 +180,7 @@ gpu_aamp_stimp.__doc__ = ast.get_docstring(cd) # Fix _gpu_searchsorted_left Docs - _gpu_searchsorted_left.__doc__ = "" + core._gpu_searchsorted_left.__doc__ = "" filepath = pathlib.Path(__file__).parent / "core.py" file_contents = "" @@ -191,10 +192,10 @@ ] for fd in function_definitions: if fd.name == "_gpu_searchsorted_left": - _gpu_searchsorted_left.__doc__ = ast.get_docstring(fd) + core._gpu_searchsorted_left.__doc__ = ast.get_docstring(fd) # Fix _gpu_searchsorted_right Docs - _gpu_searchsorted_right.__doc__ = "" + core._gpu_searchsorted_right.__doc__ = "" filepath = pathlib.Path(__file__).parent / "core.py" file_contents = "" @@ -206,7 +207,7 @@ ] for fd in function_definitions: if fd.name == "_gpu_searchsorted_right": - _gpu_searchsorted_right.__doc__ = ast.get_docstring(fd) + core._gpu_searchsorted_right.__doc__ = ast.get_docstring(fd) try: _dist = get_distribution("stumpy") From 9e0c143d7271026e17f0f8abef7516a69f8064a6 Mon Sep 17 00:00:00 2001 From: SolidAhmad Date: Sun, 4 Dec 2022 14:40:57 -0700 Subject: [PATCH 6/9] Avoid replacing docstrings for private functions --- stumpy/__init__.py | 30 ------------------------------ 1 file changed, 30 deletions(-) diff --git a/stumpy/__init__.py b/stumpy/__init__.py index 56e8974d2..1c0481487 100644 --- a/stumpy/__init__.py +++ b/stumpy/__init__.py @@ -179,36 +179,6 @@ if cd.name == "gpu_aamp_stimp": gpu_aamp_stimp.__doc__ = ast.get_docstring(cd) - # Fix _gpu_searchsorted_left Docs - core._gpu_searchsorted_left.__doc__ = "" - filepath = pathlib.Path(__file__).parent / "core.py" - - file_contents = "" - with open(filepath, encoding="utf8") as f: - file_contents = f.read() - module = ast.parse(file_contents) - function_definitions = [ - node for node in module.body if isinstance(node, ast.FunctionDef) - ] - for fd in function_definitions: - if fd.name == "_gpu_searchsorted_left": - core._gpu_searchsorted_left.__doc__ = ast.get_docstring(fd) - - # Fix _gpu_searchsorted_right Docs - core._gpu_searchsorted_right.__doc__ = "" - filepath = pathlib.Path(__file__).parent / "core.py" - - file_contents = "" - with open(filepath, encoding="utf8") as f: - file_contents = f.read() - module = ast.parse(file_contents) - function_definitions = [ - node for node in module.body if isinstance(node, ast.FunctionDef) - ] - for fd in function_definitions: - if fd.name == "_gpu_searchsorted_right": - core._gpu_searchsorted_right.__doc__ = ast.get_docstring(fd) - try: _dist = get_distribution("stumpy") # Normalize case for Windows systems From 0b0cf75d05b98eb0f11ceeef446245441454ee34 Mon Sep 17 00:00:00 2001 From: SolidAhmad Date: Sun, 4 Dec 2022 20:53:33 -0700 Subject: [PATCH 7/9] minor changes --- stumpy/__init__.py | 3 ++- stumpy/gpu_aamp.py | 3 +-- stumpy/gpu_stump.py | 3 +-- 3 files changed, 4 insertions(+), 5 deletions(-) diff --git a/stumpy/__init__.py b/stumpy/__init__.py index 1c0481487..e59444d20 100644 --- a/stumpy/__init__.py +++ b/stumpy/__init__.py @@ -39,7 +39,6 @@ from .gpu_stimp import gpu_stimp # noqa: F401 from .gpu_aamp_stimp import gpu_aamp_stimp # noqa: F401 else: # pragma: no cover - from . import core from .core import _gpu_stump_driver_not_found as gpu_stump # noqa: F401 from .core import _gpu_aamp_driver_not_found as gpu_aamp # noqa: F401 from .core import _gpu_ostinato_driver_not_found as gpu_ostinato # noqa: F401 @@ -51,6 +50,8 @@ from .core import _gpu_stimp_driver_not_found as gpu_stimp # noqa: F401 from .core import _gpu_aamp_stimp_driver_not_found as gpu_aamp_stimp # noqa: F401 + from . import core + core._gpu_searchsorted_left = core._gpu_searchsorted_left_driver_not_found core._gpu_searchsorted_right = core._gpu_searchsorted_right_driver_not_found diff --git a/stumpy/gpu_aamp.py b/stumpy/gpu_aamp.py index 0843a7a3f..7da2e4a9b 100644 --- a/stumpy/gpu_aamp.py +++ b/stumpy/gpu_aamp.py @@ -10,7 +10,6 @@ from numba import cuda from . import core, config -from .core import _gpu_searchsorted_right @cuda.jit( @@ -185,7 +184,7 @@ def _compute_and_update_PI_kernel( indices_R[j] = i if p_norm < profile[j, -1]: - idx = _gpu_searchsorted_right(profile[j], p_norm, bfs, nlevel) + idx = core._gpu_searchsorted_right(profile[j], p_norm, bfs, nlevel) for g in range(k - 1, idx, -1): profile[j, g] = profile[j, g - 1] indices[j, g] = indices[j, g - 1] diff --git a/stumpy/gpu_stump.py b/stumpy/gpu_stump.py index 9d6e4b19a..fa27b0bf3 100644 --- a/stumpy/gpu_stump.py +++ b/stumpy/gpu_stump.py @@ -10,7 +10,6 @@ from numba import cuda from . import core, config -from .core import _gpu_searchsorted_right from .gpu_aamp import gpu_aamp @@ -193,7 +192,7 @@ def _compute_and_update_PI_kernel( indices_R[j] = i if p_norm < profile[j, -1]: - idx = _gpu_searchsorted_right(profile[j], p_norm, bfs, nlevel) + idx = core._gpu_searchsorted_right(profile[j], p_norm, bfs, nlevel) for g in range(k - 1, idx, -1): profile[j, g] = profile[j, g - 1] indices[j, g] = indices[j, g - 1] From d3cf438ddf5939bb633fe4c1f38de6f8e856ad8e Mon Sep 17 00:00:00 2001 From: ninimama Date: Mon, 5 Dec 2022 08:56:21 -0700 Subject: [PATCH 8/9] Remove unnecessary import of function --- tests/test_core.py | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/tests/test_core.py b/tests/test_core.py index a1a2ddd78..b9e299fd6 100644 --- a/tests/test_core.py +++ b/tests/test_core.py @@ -4,7 +4,6 @@ import pandas as pd from scipy.spatial.distance import cdist from stumpy import core, config -from stumpy.core import _gpu_searchsorted_left, _gpu_searchsorted_right import pytest from unittest.mock import patch import os @@ -20,9 +19,9 @@ def _gpu_searchsorted_kernel(a, v, bfs, nlevel, is_left, idx): i = cuda.grid(1) if i < a.shape[0]: if is_left: - idx[i] = _gpu_searchsorted_left(a[i], v[i], bfs, nlevel) + idx[i] = core._gpu_searchsorted_left(a[i], v[i], bfs, nlevel) else: - idx[i] = _gpu_searchsorted_right(a[i], v[i], bfs, nlevel) + idx[i] = core._gpu_searchsorted_right(a[i], v[i], bfs, nlevel) try: From 2d2f5a60fb0f6429ccc2ebd10c7a9f20c8608bed Mon Sep 17 00:00:00 2001 From: ninimama Date: Mon, 5 Dec 2022 15:12:33 -0700 Subject: [PATCH 9/9] Removed hardcoded value for number of threads per block in GPU --- tests/test_core.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/tests/test_core.py b/tests/test_core.py index b9e299fd6..da1073099 100644 --- a/tests/test_core.py +++ b/tests/test_core.py @@ -29,6 +29,8 @@ def _gpu_searchsorted_kernel(a, v, bfs, nlevel, is_left, idx): except ModuleNotFoundError: from numba.core.errors import NumbaPerformanceWarning +TEST_THREADS_PER_BLOCK = 10 + def naive_rolling_window_dot_product(Q, T): window = len(Q) @@ -1387,7 +1389,7 @@ def test_find_matches_maxmatch(): @pytest.mark.filterwarnings("ignore", category=NumbaPerformanceWarning) -@patch("stumpy.config.STUMPY_THREADS_PER_BLOCK", 10) +@patch("stumpy.config.STUMPY_THREADS_PER_BLOCK", TEST_THREADS_PER_BLOCK) def test_gpu_searchsorted(): if not cuda.is_available(): # pragma: no cover pytest.skip("Skipping Tests No GPUs Available")