From 8881797c991314fa06cb4edbbbc71be4e80aee8a Mon Sep 17 00:00:00 2001 From: brandon-b-miller Date: Mon, 26 Sep 2022 15:29:42 -0700 Subject: [PATCH 1/2] remove strings_udf.ENABLED --- python/cudf/cudf/core/udf/__init__.py | 64 ++++++++----------- python/strings_udf/strings_udf/__init__.py | 62 ------------------ .../strings_udf/tests/test_string_udfs.py | 3 - 3 files changed, 28 insertions(+), 101 deletions(-) diff --git a/python/cudf/cudf/core/udf/__init__.py b/python/cudf/cudf/core/udf/__init__.py index c128bc2436c..33e5fd9846c 100644 --- a/python/cudf/cudf/core/udf/__init__.py +++ b/python/cudf/cudf/core/udf/__init__.py @@ -1,10 +1,6 @@ # Copyright (c) 2022, NVIDIA CORPORATION. -import numpy as np -from numba import cuda, types -from numba.cuda.cudaimpl import ( - lower as cuda_lower, - registry as cuda_lowering_registry, -) +from numba import types +from numba.cuda.cudaimpl import lower as cuda_lower from cudf.core.dtypes import dtype from cudf.core.udf import api, row_function, utils @@ -28,36 +24,32 @@ _STRING_UDFS_ENABLED = False try: import strings_udf - - if strings_udf.ENABLED: - from . import strings_typing # isort: skip - from . import strings_lowering # isort: skip - from strings_udf import ptxpath - from strings_udf._lib.cudf_jit_udf import to_string_view_array - from strings_udf._typing import str_view_arg_handler, string_view - - # add an overload of MaskedType.__init__(string_view, bool) - cuda_lower(api.Masked, strings_typing.string_view, types.boolean)( - masked_lowering.masked_constructor - ) - - # add an overload of pack_return(string_view) - cuda_lower(api.pack_return, strings_typing.string_view)( - masked_lowering.pack_return_scalar_impl - ) - - _supported_masked_types |= {strings_typing.string_view} - utils.launch_arg_getters[dtype("O")] = to_string_view_array - utils.masked_array_types[dtype("O")] = string_view - utils.JIT_SUPPORTED_TYPES |= STRING_TYPES - utils.ptx_files.append(ptxpath) - utils.arg_handlers.append(str_view_arg_handler) - row_function.itemsizes[dtype("O")] = string_view.size_bytes - - _STRING_UDFS_ENABLED = True - else: - del strings_udf - + from strings_udf import ptxpath + from strings_udf._lib.cudf_jit_udf import to_string_view_array + from strings_udf._typing import str_view_arg_handler, string_view + + from . import strings_typing # isort: skip + from . import strings_lowering # isort: skip + + # add an overload of MaskedType.__init__(string_view, bool) + cuda_lower(api.Masked, strings_typing.string_view, types.boolean)( + masked_lowering.masked_constructor + ) + + # add an overload of pack_return(string_view) + cuda_lower(api.pack_return, strings_typing.string_view)( + masked_lowering.pack_return_scalar_impl + ) + + _supported_masked_types |= {strings_typing.string_view} + utils.launch_arg_getters[dtype("O")] = to_string_view_array + utils.masked_array_types[dtype("O")] = string_view + utils.JIT_SUPPORTED_TYPES |= STRING_TYPES + utils.ptx_files.append(ptxpath) + utils.arg_handlers.append(str_view_arg_handler) + row_function.itemsizes[dtype("O")] = string_view.size_bytes + + _STRING_UDFS_ENABLED = True except ImportError as e: # allow cuDF to work without strings_udf pass diff --git a/python/strings_udf/strings_udf/__init__.py b/python/strings_udf/strings_udf/__init__.py index 94bd2531779..033642fc1d8 100644 --- a/python/strings_udf/strings_udf/__init__.py +++ b/python/strings_udf/strings_udf/__init__.py @@ -10,66 +10,4 @@ from . import _version -ENABLED = False - - -def compiler_from_ptx_file(path): - """Parse a PTX file header and extract the CUDA version used to compile it. - - Here is an example PTX header that this function should parse: - - // Generated by NVIDIA NVVM Compiler - // - // Compiler Build ID: CL-30672275 - // Cuda compilation tools, release 11.5, V11.5.119 - // Based on NVVM 7 - """ - file = open(path).read() - major, minor = ( - re.search(r"Cuda compilation tools, release ([0-9\.]+)", file) - .group(1) - .split(".") - ) - return int(major), int(minor) - - -# adapted from PTXCompiler -cp = subprocess.run([sys.executable, "-c", CMD], capture_output=True) -if cp.returncode == 0: - # must have a driver to proceed - versions = [int(s) for s in cp.stdout.strip().split()] - driver_version = tuple(versions[:2]) - runtime_version = tuple(versions[2:]) - - # CUDA enhanced compatibility not yet enabled - if driver_version >= runtime_version: - # Load the highest compute capability file available that is less than - # the current device's. - files = glob.glob( - os.path.join(os.path.dirname(__file__), "shim_*.ptx") - ) - dev = cuda.get_current_device() - cc = "".join(str(x) for x in dev.compute_capability) - files = glob.glob( - os.path.join(os.path.dirname(__file__), "shim_*.ptx") - ) - if len(files) == 0: - raise RuntimeError( - "This strings_udf installation is missing the necessary PTX " - "files. Please file an issue reporting this error and how you " - "installed cudf and strings_udf." - ) - sms = [ - os.path.basename(f).rstrip(".ptx").lstrip("shim_") for f in files - ] - selected_sm = max(sm for sm in sms if sm < cc) - ptxpath = os.path.join( - os.path.dirname(__file__), f"shim_{selected_sm}.ptx" - ) - - if driver_version >= compiler_from_ptx_file(ptxpath): - ENABLED = True - else: - del ptxpath - __version__ = _version.get_versions()["version"] diff --git a/python/strings_udf/strings_udf/tests/test_string_udfs.py b/python/strings_udf/strings_udf/tests/test_string_udfs.py index 9038f4cc79a..d582208e5d6 100644 --- a/python/strings_udf/strings_udf/tests/test_string_udfs.py +++ b/python/strings_udf/strings_udf/tests/test_string_udfs.py @@ -15,9 +15,6 @@ from strings_udf._lib.cudf_jit_udf import to_string_view_array from strings_udf._typing import str_view_arg_handler, string_view -if not strings_udf.ENABLED: - pytest.skip("Strings UDF not enabled.", allow_module_level=True) - def get_kernel(func, dtype): """ From 73746555dc898ca0fbc1e4ccc4c02f1338f2aa7d Mon Sep 17 00:00:00 2001 From: brandon-b-miller Date: Mon, 26 Sep 2022 15:32:54 -0700 Subject: [PATCH 2/2] update build script --- ci/gpu/build.sh | 23 +++++------------------ 1 file changed, 5 insertions(+), 18 deletions(-) diff --git a/ci/gpu/build.sh b/ci/gpu/build.sh index 118bdb263af..f6db04fd02f 100755 --- a/ci/gpu/build.sh +++ b/ci/gpu/build.sh @@ -266,27 +266,14 @@ py.test -n 8 --cache-clear --basetemp="$WORKSPACE/custreamz-cuda-tmp" --junitxml gpuci_logger "Installing strings_udf" gpuci_mamba_retry install strings_udf -c "${CONDA_BLD_DIR}" -c "${CONDA_ARTIFACT_PATH}" +# only install strings_udf after cuDF is finished testing without its presence cd "$WORKSPACE/python/strings_udf/strings_udf" gpuci_logger "Python py.test for strings_udf" -# We do not want to exit with a nonzero exit code in the case where no -# strings_udf tests are run because that will always happen when the local CUDA -# version is not 11.5. We need to suppress the exit code because this script is -# run with set -e and we're already setting a trap that we don't want to -# override here. - -STRINGS_UDF_PYTEST_RETCODE=0 -py.test -n 8 --cache-clear --basetemp="$WORKSPACE/strings-udf-cuda-tmp" --junitxml="$WORKSPACE/junit-strings-udf.xml" -v --cov-config=.coveragerc --cov=strings_udf --cov-report=xml:"$WORKSPACE/python/strings_udf/strings-udf-coverage.xml" --cov-report term tests || STRINGS_UDF_PYTEST_RETCODE=$? - -if [ ${STRINGS_UDF_PYTEST_RETCODE} -eq 5 ]; then - echo "No strings UDF tests were run, but this script will continue to execute." -elif [ ${STRINGS_UDF_PYTEST_RETCODE} -ne 0 ]; then - exit ${STRINGS_UDF_PYTEST_RETCODE} -else - cd "$WORKSPACE/python/cudf/cudf" - gpuci_logger "Python py.test retest cuDF UDFs" - py.test tests/test_udf_masked_ops.py -n 8 --cache-clear -fi +# retest cudf with strings_udf present +py.test -n 8 --cache-clear --basetemp="$WORKSPACE/strings-udf-cuda-tmp" --junitxml="$WORKSPACE/junit-strings-udf.xml" -v --cov-config=.coveragerc --cov=strings_udf --cov-report=xml:"$WORKSPACE/python/strings_udf/strings-udf-coverage.xml" --cov-report term tests +gpuci_logger "Python py.test retest cuDF UDFs" +py.test tests/test_udf_masked_ops.py -n 8 --cache-clear # Run benchmarks with both cudf and pandas to ensure compatibility is maintained. # Benchmarks are run in DEBUG_ONLY mode, meaning that only small data sizes are used.