From 54b6e3f87a58ba7f1da138e729c4d312c531d5c1 Mon Sep 17 00:00:00 2001 From: Jim Crist-Harif Date: Fri, 15 May 2026 11:58:42 -0500 Subject: [PATCH 01/17] Update `libcuml` wheel to depend on `librmm` (#8110) This is required to import `libcuml`, but wasn't explicitly listed as a dependency. The parent main user-facing package (`cuml`) does have a `rmm` dependency, which in turn depends on `librmm`, so everything would be installed properly normally anyway. Just adding an explicit link here for tidyness. Part of #7845. Authors: - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Kyle Edwards (https://github.com/KyleFromNVIDIA) URL: https://github.com/rapidsai/cuml/pull/8110 --- dependencies.yaml | 1 + python/libcuml/pyproject.toml | 1 + 2 files changed, 2 insertions(+) diff --git a/dependencies.yaml b/dependencies.yaml index f1d04fec4a..2f5ceb9f9c 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -248,6 +248,7 @@ files: - cuda_wheels - depends_on_libcuvs - depends_on_libraft + - depends_on_librmm - depends_on_rapids_logger channels: - rapidsai-nightly diff --git a/python/libcuml/pyproject.toml b/python/libcuml/pyproject.toml index cc7a86a1ab..523df2d06f 100644 --- a/python/libcuml/pyproject.toml +++ b/python/libcuml/pyproject.toml @@ -27,6 +27,7 @@ classifiers = [ dependencies = [ "cuda-toolkit[cublas,cufft,curand,cusolver,cusparse]==13.*", "libraft==26.6.*,>=0.0.0a0", + "librmm==26.6.*,>=0.0.0a0", "nvidia-nvjitlink>=13.0,<14", "rapids-logger==0.2.*,>=0.0.0a0", ] # This list was generated by `rapids-dependency-file-generator`. To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. From da8e0cd567c98a831968ff8142b9409e0abc3931 Mon Sep 17 00:00:00 2001 From: Jim Crist-Harif Date: Fri, 15 May 2026 13:36:50 -0500 Subject: [PATCH 02/17] Fix wrapping metaestimators in `Pipeline` in `cuml.accel` (#8115) Our pipeline data transfer optimization didn't work if any of the steps were other compositional metaestimators that wrapped accelerated estimators (since these could then accidentally use the accelerated versions, resulting in a mix of `cupy` and `numpy` results). This PR patches the other two compositional estimators (`FeatureUnion` and `ColumnTransformer`) so they always run within a `numpy` output-type context. Fixes #8112. Fixes a few sklearn examples as well (yay!) Authors: - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Simon Adorf (https://github.com/csadorf) URL: https://github.com/rapidsai/cuml/pull/8115 --- .../cuml/accel/_patches/sklearn/compose.py | 25 +++++++++ .../cuml/accel/_patches/sklearn/pipeline.py | 24 ++++++-- python/cuml/cuml/accel/core.py | 1 + python/cuml/cuml_accel_tests/test_pipeline.py | 55 ++++++++++++++++++- .../upstream/scikit-learn/xfail-examples.yaml | 5 -- 5 files changed, 99 insertions(+), 11 deletions(-) create mode 100644 python/cuml/cuml/accel/_patches/sklearn/compose.py diff --git a/python/cuml/cuml/accel/_patches/sklearn/compose.py b/python/cuml/cuml/accel/_patches/sklearn/compose.py new file mode 100644 index 0000000000..8f8f9dec54 --- /dev/null +++ b/python/cuml/cuml/accel/_patches/sklearn/compose.py @@ -0,0 +1,25 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. +# SPDX-License-Identifier: Apache-2.0 +import functools + +from sklearn.compose import ColumnTransformer + +from cuml.internals.outputs import using_output_type + +__all__ = ("ColumnTransformer",) + + +def patch_method(name): + """Patch a ColumnTransformer method to ensure results returned as numpy.""" + orig_method = getattr(ColumnTransformer, name) + + @functools.wraps(orig_method) + def method(self, *args, **kwargs): + with using_output_type("numpy"): + return orig_method(self, *args, **kwargs) + + setattr(ColumnTransformer, name, method) + + +for method_name in ["fit", "fit_transform", "transform"]: + patch_method(method_name) diff --git a/python/cuml/cuml/accel/_patches/sklearn/pipeline.py b/python/cuml/cuml/accel/_patches/sklearn/pipeline.py index 629f0ea86a..a032f290a0 100644 --- a/python/cuml/cuml/accel/_patches/sklearn/pipeline.py +++ b/python/cuml/cuml/accel/_patches/sklearn/pipeline.py @@ -5,14 +5,14 @@ import cupy as cp from cupyx.scipy.sparse import issparse as is_cp_sparse -from sklearn.pipeline import Pipeline +from sklearn.pipeline import FeatureUnion, Pipeline from sklearn.utils.metaestimators import available_if from cuml.accel.estimator_proxy import is_proxy from cuml.internals.global_settings import GlobalSettings from cuml.internals.outputs import using_output_type -__all__ = ("Pipeline",) +__all__ = ("Pipeline", "FeatureUnion") def get_output_type(pipeline, reverse=False): @@ -54,7 +54,7 @@ def flat_steps(pipeline): return "cupy" -def patch_method(name): +def patch_pipeline_method(name): """Patch a sklearn Pipeline method to reduce device<->host transfers.""" orig_method = inspect.getattr_static(Pipeline, name) # Unwrap @available_if decorated methods @@ -105,4 +105,20 @@ def method(self, *args, **kwargs): "score_samples", "transform", ]: - patch_method(method_name) + patch_pipeline_method(method_name) + + +def patch_feature_union_method(name): + """Patch a FeatureUnion method to ensure results returned as numpy.""" + orig_method = getattr(FeatureUnion, name) + + @functools.wraps(orig_method) + def method(self, *args, **kwargs): + with using_output_type("numpy"): + return orig_method(self, *args, **kwargs) + + setattr(FeatureUnion, name, method) + + +for method_name in ["fit", "fit_transform", "transform"]: + patch_feature_union_method(method_name) diff --git a/python/cuml/cuml/accel/core.py b/python/cuml/cuml/accel/core.py index 22d15868e4..8ea280d054 100644 --- a/python/cuml/cuml/accel/core.py +++ b/python/cuml/cuml/accel/core.py @@ -92,6 +92,7 @@ def debug(self, msg: str) -> None: _PATCHES = { "sklearn.pipeline", + "sklearn.compose", "sklearn.utils", "sklearn.utils._array_api", "sklearn.utils.discovery", diff --git a/python/cuml/cuml_accel_tests/test_pipeline.py b/python/cuml/cuml_accel_tests/test_pipeline.py index cb0935d554..814b114ad2 100644 --- a/python/cuml/cuml_accel_tests/test_pipeline.py +++ b/python/cuml/cuml_accel_tests/test_pipeline.py @@ -13,6 +13,7 @@ from packaging.version import Version from sklearn.base import BaseEstimator from sklearn.cluster import DBSCAN, KMeans +from sklearn.compose import ColumnTransformer from sklearn.datasets import make_classification, make_regression from sklearn.decomposition import PCA, TruncatedSVD from sklearn.linear_model import ( @@ -28,8 +29,8 @@ KNeighborsRegressor, NearestNeighbors, ) -from sklearn.pipeline import Pipeline, make_pipeline -from sklearn.preprocessing import StandardScaler +from sklearn.pipeline import FeatureUnion, Pipeline, make_pipeline +from sklearn.preprocessing import RobustScaler, StandardScaler from umap import UMAP SKLEARN_18 = Version(sklearn.__version__) >= Version("1.8.0.dev0") @@ -356,3 +357,53 @@ def test_pipeline_classifier_predict_non_numeric_labels(patch_methods): assert isinstance(LogisticRegression.predict.args[0], cp.ndarray) # User-facing output is always numpy assert isinstance(out, np.ndarray) + + +@requires_sklearn_18 +def test_column_transfomer_in_pipeline_works(): + """Ensure outputs of steps in `ColumnTransformer` return as numpy""" + rng = np.random.default_rng(0) + X = rng.standard_normal((200, 20)).astype(np.float32) + y = rng.standard_normal(200).astype(np.float32) + + ct = ColumnTransformer( + [ + ("svd", TruncatedSVD(n_components=5), slice(0, 10)), + ("pass", "passthrough", slice(10, 20)), + ] + ) + + pipe = Pipeline( + [ + ("ct", ct), # Shouldn't be accelerated + ("scaler", RobustScaler()), # Not accelerated + ("ridge", Ridge()), # Accelerated + ] + ) + + pipe.fit(X, y) + + +@requires_sklearn_18 +def test_feature_union_in_pipeline_works(): + """Ensure outputs of steps in `FeatureUnion` return as numpy""" + rng = np.random.default_rng(0) + X = rng.standard_normal((200, 20)).astype(np.float32) + y = rng.standard_normal(200).astype(np.float32) + + union = FeatureUnion( + [ + ("svd", TruncatedSVD(n_components=2)), + ("pca", PCA(n_components=2)), + ] + ) + + pipe = Pipeline( + [ + ("features", union), # Shouldn't be accelerated + ("scaler", RobustScaler()), # Not accelerated + ("ridge", Ridge()), # Accelerated + ] + ) + + pipe.fit(X, y) diff --git a/python/cuml/cuml_accel_tests/upstream/scikit-learn/xfail-examples.yaml b/python/cuml/cuml_accel_tests/upstream/scikit-learn/xfail-examples.yaml index e19cee5cbf..0027602b97 100644 --- a/python/cuml/cuml_accel_tests/upstream/scikit-learn/xfail-examples.yaml +++ b/python/cuml/cuml_accel_tests/upstream/scikit-learn/xfail-examples.yaml @@ -3,11 +3,6 @@ strict: false tests: - "decomposition::plot_faces_decomposition" -- reason: 'cuml.accel bug: implicit CuPy-to-NumPy conversion not handled in ColumnTransformer/Pipeline' - marker: cuml_accel_bugs - tests: - - "linear_model::plot_poisson_regression_non_normal_loss" - - "release_highlights::plot_release_highlights_1_1_0" - reason: 'cuml.accel bug: native crash in cuml PCA' marker: cuml_accel_bugs tests: From 0095e1e2edf2ac9be9a30185be155ca80f231504 Mon Sep 17 00:00:00 2001 From: Jim Crist-Harif Date: Fri, 15 May 2026 14:32:21 -0500 Subject: [PATCH 03/17] Use the PTDS for (most) cupy operations (#8086) This configures cupy to use the per-thread default stream (PTDS) for _most_ operations. This avoids usage of the default legacy stream in more of the codebase, allowing for improved parallelism when running across multiple threads. **This is a breaking change.** Previously any cupy operations in `cuml` ran in cupy's default stream (the legacy stream). We didn't synchronize the stream before returning, but that didn't matter due to the synchronization behavior of the legacy stream. With this PR we've moved to running (most) cupy operations in the PTDS. Depending on the operation, we may not synchronize the PTDS before returning. **Most users shouldn't notice a difference and should have no issues.** Users not using threads, custom streams, or only working with host memory (e.g. numpy in/numpy out) should see no difference. Likewise any users that only use cupy's default stream (the legacy stream) in their code should see no issues. Users doing tricky things with custom streams or threads may run into issues and require a manual sync of the PTDS (can be done with `cupy.cuda.Stream.ptds.synchronize()`. For example, the following workflow _may_ run into issues: - Run a cuml operation based on cupy in thread A, returning a cupy array - Consume that output in thread B as a cupy array using a stream other than the legacy stream (e.g. a different PTDS or a custom stream) For safety, you probably want to add a call to `cupy.cuda.Stream.ptds.synchronize()` in thread A before returning to ensure the output array is fully populated before consuming it in thread B. Fixes #7909. Authors: - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Simon Adorf (https://github.com/csadorf) - Dante Gama Dessavre (https://github.com/dantegd) URL: https://github.com/rapidsai/cuml/pull/8086 --- python/cuml/cuml/internals/outputs.py | 7 +++-- python/cuml/cuml/preprocessing/label.py | 4 --- python/cuml/tests/test_prims.py | 8 ----- python/cuml/tests/test_reflection.py | 40 +++++++++++++++++++++++++ 4 files changed, 45 insertions(+), 14 deletions(-) diff --git a/python/cuml/cuml/internals/outputs.py b/python/cuml/cuml/internals/outputs.py index 2b91978f81..0bcdb535c3 100644 --- a/python/cuml/cuml/internals/outputs.py +++ b/python/cuml/cuml/internals/outputs.py @@ -7,6 +7,7 @@ import inspect import numpy as np +from cupy.cuda import Stream # TODO: Try to resolve circular import that makes this necessary: from cuml.internals import input_utils as iu @@ -212,7 +213,8 @@ def enter_internal_context(): gs._external_output_type = gs.output_type gs.output_type = "mirror" try: - yield True + with Stream.ptds: + yield True finally: gs.output_type = gs._external_output_type gs._external_output_type = False @@ -451,6 +453,7 @@ def inner(*args, **kwargs): # We're internal, return as cuml output_type = "cuml" - return coerce_arrays(res, output_type) + with enter_internal_context(): + return coerce_arrays(res, output_type) return inner diff --git a/python/cuml/cuml/preprocessing/label.py b/python/cuml/cuml/preprocessing/label.py index 590e5fc458..a042130c8e 100644 --- a/python/cuml/cuml/preprocessing/label.py +++ b/python/cuml/cuml/preprocessing/label.py @@ -49,8 +49,6 @@ def label_binarize( dtype=cp.float32, ) - cp.cuda.Stream.null.synchronize() - is_binary = classes.shape[0] == 2 if sparse_output: @@ -185,8 +183,6 @@ def fit(self, y) -> "LabelBinarizer": else: self.classes_ = cp.unique(y).astype(y.dtype) - cp.cuda.Stream.null.synchronize() - return self @cuml.internals.reflect diff --git a/python/cuml/tests/test_prims.py b/python/cuml/tests/test_prims.py index 30e4fcccca..5f497bb279 100644 --- a/python/cuml/tests/test_prims.py +++ b/python/cuml/tests/test_prims.py @@ -24,8 +24,6 @@ def test_monotonic_without_classes(arr_type, dtype, copy): monotonic, returned_classes = make_monotonic(arr, copy=copy) - cp.cuda.Stream.null.synchronize() - # Verify monotonic mapping: [0, 15, 10, 50, 20, 50] -> [0, 2, 1, 4, 3, 4] # (sorted unique: 0->0, 10->1, 15->2, 20->3, 50->4) expected_monotonic = cp.array([0, 2, 1, 4, 3, 4], dtype=dtype) @@ -56,8 +54,6 @@ def test_monotonic_inversion(dtype): # Invert: use classes array to map indices back to original values inverted = classes[monotonic] - cp.cuda.Stream.null.synchronize() - assert array_equal(inverted, original) @@ -74,8 +70,6 @@ def test_monotonic_with_explicit_classes(dtype, copy): labels, classes=classes, copy=copy ) - cp.cuda.Stream.null.synchronize() - # Labels should map to their position in the original classes array # 5 -> 0, 2 -> 1, 8 -> 2 expected = cp.array([2, 1, 0, 1, 2], dtype=dtype) @@ -100,8 +94,6 @@ def test_monotonic_unknown_labels(dtype): monotonic, _ = make_monotonic(labels, classes=classes, copy=True) - cp.cuda.Stream.null.synchronize() - # Unknown labels (999, -1) should map to len(classes) = 3 # 1 -> 0, 999 -> 3, 2 -> 1, 3 -> 2, -1 -> 3 expected = cp.array([0, 3, 1, 2, 3], dtype=dtype) diff --git a/python/cuml/tests/test_reflection.py b/python/cuml/tests/test_reflection.py index 9931e4e432..08b2560400 100644 --- a/python/cuml/tests/test_reflection.py +++ b/python/cuml/tests/test_reflection.py @@ -394,3 +394,43 @@ def test_array_descriptor_cache_behavior(): assert b"pandas" not in msg assert_output_type(model2.X_, "cupy") assert len(model2.__dict__["X_"].values) == 2 # cuml + cupy + + +def test_decorators_set_cupy_ptds(): + class MyEstimator(Base): + @reflect(reset="type") + def fit(self, X, y=None): + assert cp.cuda.get_current_stream() is cp.cuda.Stream.ptds + return self + + @reflect + def direct_call(self, X): + assert cp.cuda.get_current_stream() is cp.cuda.Stream.ptds + return cp.zeros(3) + + @reflect + def nested_call(self, X): + assert cp.cuda.get_current_stream() is cp.cuda.Stream.ptds + return self.direct_call(X) + + @run_in_internal_context + def no_reflection(self, X): + assert cp.cuda.get_current_stream() is cp.cuda.Stream.ptds + return cp.zeros(3) + + X = cp.ones(3) + + # Check that ptds is used instead of the default stream + model = MyEstimator() + model.fit(X) + model.direct_call(X) + model.nested_call(X) + model.no_reflection(X) + + # Check that ptds is used instead of a custom stream + with cp.cuda.Stream(): + model = MyEstimator() + model.fit(X) + model.direct_call(X) + model.nested_call(X) + model.no_reflection(X) From 4ad59e410a6908bb4ca8db2f64043b07fc1b49e7 Mon Sep 17 00:00:00 2001 From: Tim Head Date: Fri, 15 May 2026 21:33:34 +0200 Subject: [PATCH 04/17] DOC Add third-party app example for cuml.accel (#8094) This adds to the cuml.accel documentation and documents the "third party application" use-case. Authors: - Tim Head (https://github.com/betatim) Approvers: - Jim Crist-Harif (https://github.com/jcrist) URL: https://github.com/rapidsai/cuml/pull/8094 --- docs/source/cuml-accel/examples/index.rst | 1 + .../cuml-accel/examples/third-party-apps.rst | 117 ++++++++++++++++++ docs/source/cuml-accel/index.rst | 6 + 3 files changed, 124 insertions(+) create mode 100644 docs/source/cuml-accel/examples/third-party-apps.rst diff --git a/docs/source/cuml-accel/examples/index.rst b/docs/source/cuml-accel/examples/index.rst index de65734403..3ea9b53186 100644 --- a/docs/source/cuml-accel/examples/index.rst +++ b/docs/source/cuml-accel/examples/index.rst @@ -11,4 +11,5 @@ examples in this section is available in the cuML GitHub repository at `examples getting_started.ipynb profiling.ipynb plot_kmeans_digits.ipynb + third-party-apps.rst onnx_export.ipynb diff --git a/docs/source/cuml-accel/examples/third-party-apps.rst b/docs/source/cuml-accel/examples/third-party-apps.rst new file mode 100644 index 0000000000..85449af0ee --- /dev/null +++ b/docs/source/cuml-accel/examples/third-party-apps.rst @@ -0,0 +1,117 @@ +Accelerating Third-Party Applications +====================================== + +The ``CUML_ACCEL_ENABLED`` environment variable lets you GPU-accelerate any +Python application that uses ``sklearn``, ``umap``, or ``hdbscan``. +Even applications whose code you cannot modify. This is useful for +installed CLI tools, applications, and third-party libraries. + +.. code-block:: console + + CUML_ACCEL_ENABLED=1 some-third-party-tool [args...] + +When :ref:`CUML_ACCEL_ENABLED=1 is defined `, +`cuml.accel` will be enabled as part of the normal Python interpreter +startup, letting you accelerate Python applications without modification + +This means you do not need access to an application's source code: set the +environment variable and the acceleration applies automatically. + +Example: Embedding Visualization with embedding-atlas +----------------------------------------------------- + +`embedding-atlas `_ is Apple's +open-source tool for interactive visualization of large embedding datasets. +Given a text dataset, it computes sentence embeddings, projects them to 2D +using `UMAP `_, and launches a +browser-based explorer. + +Install it alongside ``cuml``: + +.. code-block:: console + + pip install embedding-atlas + +Run it on a Hugging Face dataset. The example below uses +`TinyStories `_, +a dataset of 2M+ short stories: + +.. code-block:: console + + # CPU -- UMAP runs on CPU + embedding-atlas roneneldan/TinyStories --text text \ + --split train --sample 1000000 + + # GPU -- set environment variable; no other changes needed + CUML_ACCEL_ENABLED=1 embedding-atlas roneneldan/TinyStories --text text \ + --split train --sample 1000000 + +The only change between the two commands is the environment variable. +``embedding-atlas`` computes embeddings with sentence-transformers (which +already uses the GPU), then runs UMAP for dimensionality reduction. +``cuml.accel`` intercepts the ``umap.UMAP`` call inside ``embedding-atlas`` +and dispatches ``fit_transform`` to cuML's GPU implementation. + +Use a smaller ``--sample`` value (e.g. 250000) for a quicker test run. +The UMAP speedup grows with dataset size. + +To confirm GPU dispatch, add ``CUML_ACCEL_LOG_LEVEL=info``: + +.. code-block:: console + + CUML_ACCEL_ENABLED=1 CUML_ACCEL_LOG_LEVEL=info embedding-atlas \ + roneneldan/TinyStories --text text --split train --sample 1000000 + +You should see the following messages amongst the other output: + +.. code-block:: text + + [cuml.accel] Accelerator installed. + [cuml.accel] `UMAP.fit_transform` ran on GPU + +Results +~~~~~~~ + +At the time of writing and on the hardware the author used the +``fit_transform`` step saw a roughly **~4x speedup** because cuML's GPU +UMAP replaces the CPU optimization. The KNN step (``nearest_neighbors``) +is a standalone function call that ``cuml.accel`` does not currently +intercept, so it runs on CPU in both cases. Despite this, the overall +UMAP step is still **~2x faster**. + +At smaller scales (< 100K rows) the UMAP step is already fast on CPU and +the speedup is less pronounced. The benefit grows with dataset size. + + +Identifying Acceleratable Applications +--------------------------------------- + +Any Python tool that calls one of the following is a candidate for +``CUML_ACCEL_ENABLED``: + +- ``sklearn`` estimators (KMeans, PCA, DBSCAN, RandomForest, + LogisticRegression, NearestNeighbors, and + :doc:`many more <../faq>`) +- ``umap.UMAP`` +- ``hdbscan.HDBSCAN`` + +A quick way to check: search an application's dependencies for +``scikit-learn``, ``umap-learn``, or ``hdbscan``, or run with +``CUML_ACCEL_LOG_LEVEL=info`` and look for ``ran on GPU`` messages +in the output. + +Checking for CPU Fallbacks +-------------------------- + +Not all parameter combinations are supported on the GPU. When +``cuml.accel`` encounters an unsupported configuration, it silently +falls back to CPU execution. To detect this, set the log level to +``info`` or ``debug``: + +.. code-block:: console + + CUML_ACCEL_ENABLED=1 CUML_ACCEL_LOG_LEVEL=info python app.py + +Lines containing ``ran on GPU`` confirm GPU execution. Lines +containing ``falling back to CPU`` indicate a fallback, along with +the reason. See :doc:`../logging-and-profiling` for more detail. diff --git a/docs/source/cuml-accel/index.rst b/docs/source/cuml-accel/index.rst index 4caadb1e28..0f1ca42196 100644 --- a/docs/source/cuml-accel/index.rst +++ b/docs/source/cuml-accel/index.rst @@ -58,6 +58,8 @@ executing the following line magic at the top (before other imports): You can see an example of this in :doc:`this example `. +.. _cuml-accel-env-var: + Environment Variable ~~~~~~~~~~~~~~~~~~~~ @@ -75,6 +77,10 @@ environment variable to ``1`` or ``true`` (case insensitive). Note that any python program running with the environment defined this way will load the accelerator, which may result in a measurable startup overhead. +This approach is especially useful for accelerating +:doc:`third-party applications ` whose code you do not +control. + Additionally, if ``cuml`` is not installed properly in your environment, the ``CUML_ACCEL_ENABLED`` environment variable will be silently ignored (and normal CPU execution will occur). For this reason one of the other methods From 87d7b8ea0f099b10f72e097a8a1ab3088650d831 Mon Sep 17 00:00:00 2001 From: Jim Crist-Harif Date: Fri, 15 May 2026 15:50:26 -0500 Subject: [PATCH 05/17] Numpy 1.x compatibility fixes (#8118) This: - Bumps our minimum supported `numpy` version to 1.26, to match that of `cudf`. Since `cudf` is a required dependency, we were effectively pinned at that already. - Adds `numpy` to our oldest deps test runs. This also effectively adds `cupy==13.6`, since `cupy==14` requires `numpy>=2.0`. Explicitly specifying `cupy==13.6` in an oldest-deps run is tricky since the pypi packages require cuda suffixes as well. I'm skipping that for now. - Fixes a few incompatibilities with numpy 1.x Authors: - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Gil Forsyth (https://github.com/gforsyth) - Simon Adorf (https://github.com/csadorf) URL: https://github.com/rapidsai/cuml/pull/8118 --- .../all_cuda-129_arch-aarch64.yaml | 2 +- .../all_cuda-129_arch-x86_64.yaml | 2 +- .../all_cuda-132_arch-aarch64.yaml | 2 +- .../all_cuda-132_arch-x86_64.yaml | 2 +- conda/recipes/cuml/recipe.yaml | 2 +- dependencies.yaml | 3 +- python/cuml/cuml/internals/validation.py | 29 ++++++++++++------- .../cuml/linear_model/linear_regression.pyx | 2 +- python/cuml/cuml/linear_model/ridge.pyx | 2 +- python/cuml/pyproject.toml | 2 +- .../cuml/tests/explainer/test_gpu_treeshap.py | 6 ++++ python/cuml/tests/test_validation.py | 5 ++-- 12 files changed, 38 insertions(+), 21 deletions(-) diff --git a/conda/environments/all_cuda-129_arch-aarch64.yaml b/conda/environments/all_cuda-129_arch-aarch64.yaml index 699e9872f3..f70b17a786 100644 --- a/conda/environments/all_cuda-129_arch-aarch64.yaml +++ b/conda/environments/all_cuda-129_arch-aarch64.yaml @@ -47,7 +47,7 @@ dependencies: - nltk - numba-cuda>=0.22.2,<0.29.0 - numba>=0.60.0,<0.65.0 -- numpy>=1.23,<3.0 +- numpy>=1.26,<3.0 - numpydoc - numpydoc<1.9 - nvidia-ml-py>=12 diff --git a/conda/environments/all_cuda-129_arch-x86_64.yaml b/conda/environments/all_cuda-129_arch-x86_64.yaml index 644e0132e7..130a281624 100644 --- a/conda/environments/all_cuda-129_arch-x86_64.yaml +++ b/conda/environments/all_cuda-129_arch-x86_64.yaml @@ -46,7 +46,7 @@ dependencies: - nltk - numba-cuda>=0.22.2,<0.29.0 - numba>=0.60.0,<0.65.0 -- numpy>=1.23,<3.0 +- numpy>=1.26,<3.0 - numpydoc - numpydoc<1.9 - nvidia-ml-py>=12 diff --git a/conda/environments/all_cuda-132_arch-aarch64.yaml b/conda/environments/all_cuda-132_arch-aarch64.yaml index 692f50dbd6..40981932cf 100644 --- a/conda/environments/all_cuda-132_arch-aarch64.yaml +++ b/conda/environments/all_cuda-132_arch-aarch64.yaml @@ -47,7 +47,7 @@ dependencies: - nltk - numba-cuda>=0.22.2,<0.29.0 - numba>=0.60.0,<0.65.0 -- numpy>=1.23,<3.0 +- numpy>=1.26,<3.0 - numpydoc - numpydoc<1.9 - nvidia-ml-py>=12 diff --git a/conda/environments/all_cuda-132_arch-x86_64.yaml b/conda/environments/all_cuda-132_arch-x86_64.yaml index 2eda3604e5..b879ce0d48 100644 --- a/conda/environments/all_cuda-132_arch-x86_64.yaml +++ b/conda/environments/all_cuda-132_arch-x86_64.yaml @@ -46,7 +46,7 @@ dependencies: - nltk - numba-cuda>=0.22.2,<0.29.0 - numba>=0.60.0,<0.65.0 -- numpy>=1.23,<3.0 +- numpy>=1.26,<3.0 - numpydoc - numpydoc<1.9 - nvidia-ml-py>=12 diff --git a/conda/recipes/cuml/recipe.yaml b/conda/recipes/cuml/recipe.yaml index 2c526f86ca..960d4347df 100644 --- a/conda/recipes/cuml/recipe.yaml +++ b/conda/recipes/cuml/recipe.yaml @@ -100,7 +100,7 @@ requirements: - libcuml =${{ version }} - numba >=0.60.0,<0.65.0 - numba-cuda >=0.22.2,<0.29.0 - - numpy >=1.23,<3.0 + - numpy >=1.26,<3.0 - scikit-learn >=1.4 - scipy >=1.14.0 - packaging diff --git a/dependencies.yaml b/dependencies.yaml index 2f5ceb9f9c..12705a4401 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -324,7 +324,7 @@ dependencies: packages: - joblib>=0.11 - numba>=0.60.0,<0.65.0 - - &numpy numpy>=1.23,<3.0 + - &numpy numpy>=1.26,<3.0 - scipy>=1.14.0 - packaging - rich @@ -531,6 +531,7 @@ dependencies: - scikit-learn==1.5.0 - umap-learn==0.5.7 - hdbscan==0.8.39 + - numpy==1.26 - matrix: {dependencies: "intermediate"} packages: - scikit-learn==1.7.2 diff --git a/python/cuml/cuml/internals/validation.py b/python/cuml/cuml/internals/validation.py index 2e00a2c623..e8afe0aa55 100644 --- a/python/cuml/cuml/internals/validation.py +++ b/python/cuml/cuml/internals/validation.py @@ -351,7 +351,7 @@ def check_all_finite(array, *, allow_nan=False, input_name=None) -> None: input_name : str or None, default=None The input parameter name to use in error messages. """ - if not np.isdtype(array.dtype, "real floating"): + if not array.dtype.kind == "f": # No-op for non floating inputs return @@ -484,6 +484,19 @@ def _index_as_mem_type(index, mem_type=None): return index +if np.lib.NumpyVersion(np.__version__) >= "2.0.0b1": + np_asarray = np.asarray +else: + + def np_asarray(x, dtype=None, order=None, copy=None): + """A compatibility shim for `np.asarray`. + + numpy 2.0 added the `copy` arg to `np.asarray`, as well as changed the + meaning of copy=False to "error if a copy required" rather than "only + copy if needed" (which is now `copy=None`).""" + return np.array(x, dtype=dtype, order=order or "K", copy=bool(copy)) + + def check_array( array, *, @@ -600,7 +613,7 @@ def check_array( # Infer proper output dtype if array_dtype is not None: # Check for complex inputs before conversion when possible - if np.isdtype(array_dtype, "complex floating"): + if array_dtype.kind == "c": raise ValueError("Complex data not supported") if dtype is None: dtype = array_dtype @@ -703,7 +716,7 @@ def check_array( elif ( mem_type is None and cudf.pandas.LOADED - and np.isdtype(array.dtype, ("numeric", "bool")) + and array.dtype.kind in "iufb" ): # We treat pandas objects with supported dtypes as device # memory when running under cudf.pandas. Note that the output @@ -732,8 +745,7 @@ def check_array( array, dtype=dtype, order=order, copy=(copy or None) ) else: - # XXX: using np.array for compat with numpy < 2 - array = np.array( + array = np_asarray( array, dtype=dtype, order=order, copy=(copy or None) ) @@ -761,7 +773,7 @@ def check_array( ) # Check for complex inputs after conversion for cases when `dtype=None` - if np.isdtype(array.dtype, "complex floating"): + if array.dtype.kind == "c": raise ValueError("Complex data not supported") # Validate data meets expected value requirements @@ -1052,10 +1064,7 @@ def check_y( input_dtype = y.dtype if mem_type is None: mem_type = "host" if isinstance(y, np.ndarray) else "device" - if ( - np.isdtype(y.dtype, ("numeric", "bool")) - and return_classes is True - ): + if y.dtype.kind in "iufb" and return_classes is True: y = cp.asarray(y) elif ( y.dtype == "object" diff --git a/python/cuml/cuml/linear_model/linear_regression.pyx b/python/cuml/cuml/linear_model/linear_regression.pyx index 450db871a1..387a0dd437 100644 --- a/python/cuml/cuml/linear_model/linear_regression.pyx +++ b/python/cuml/cuml/linear_model/linear_regression.pyx @@ -241,7 +241,7 @@ class LinearRegression(Base, ) # All libcuml solvers require F-ordered X, and mutate the inputs. - X = cp.asarray(X, order="F", copy=None if may_mutate_X else True) + X = cp.array(X, order="F", copy=None if may_mutate_X else True) if not may_mutate_y: y = y.copy() if sample_weight is not None and not may_mutate_sample_weight: diff --git a/python/cuml/cuml/linear_model/ridge.pyx b/python/cuml/cuml/linear_model/ridge.pyx index c9e196cbfd..085474e848 100644 --- a/python/cuml/cuml/linear_model/ridge.pyx +++ b/python/cuml/cuml/linear_model/ridge.pyx @@ -285,7 +285,7 @@ class Ridge(Base, # The `eig` solver requires X be F-contiguous. Additionally, all inputs # are mutated when weighted or `fit_intercept=True`. mutates = self.fit_intercept or sample_weight is not None - X = cp.asarray(X, order="F", copy=True if mutates and not may_mutate_X else None) + X = cp.array(X, order="F", copy=True if mutates and not may_mutate_X else None) if mutates and not may_mutate_y: y = y.copy() if sample_weight is not None and mutates and not may_mutate_sample_weight: diff --git a/python/cuml/pyproject.toml b/python/cuml/pyproject.toml index 18a5aa8a25..04ae556bb7 100644 --- a/python/cuml/pyproject.toml +++ b/python/cuml/pyproject.toml @@ -87,7 +87,7 @@ dependencies = [ "libcuml==26.6.*,>=0.0.0a0", "numba-cuda>=0.22.2,<0.29.0", "numba>=0.60.0,<0.65.0", - "numpy>=1.23,<3.0", + "numpy>=1.26,<3.0", "nvidia-nvjitlink>=13.0,<14", "packaging", "pylibraft==26.6.*,>=0.0.0a0", diff --git a/python/cuml/tests/explainer/test_gpu_treeshap.py b/python/cuml/tests/explainer/test_gpu_treeshap.py index 3459a30ead..10b27525b7 100644 --- a/python/cuml/tests/explainer/test_gpu_treeshap.py +++ b/python/cuml/tests/explainer/test_gpu_treeshap.py @@ -124,6 +124,9 @@ def count_categorical_split(tl_model): ) def test_xgb_regressor(objective): xgb = pytest.importorskip("xgboost") + pytest.importorskip( + "numpy", minversion="2.0", reason="Test fails on numpy < 2" + ) n_samples = 100 X, y = make_regression( @@ -197,6 +200,9 @@ def test_xgb_regressor(objective): ) def test_xgb_classifier(objective, n_classes): xgb = pytest.importorskip("xgboost") + pytest.importorskip( + "numpy", minversion="2.0", reason="Test fails on numpy < 2" + ) n_samples = 100 X, y = make_classification( diff --git a/python/cuml/tests/test_validation.py b/python/cuml/tests/test_validation.py index 7d67d1c603..95a97ca113 100644 --- a/python/cuml/tests/test_validation.py +++ b/python/cuml/tests/test_validation.py @@ -841,8 +841,9 @@ def test_check_array_dataframe_mixed_dtypes(kind, mem_type): ) # Non-numeric columns -> object dtype by default if is_cuda_output(mem_type, df): - # cupy doesn't support object dtypes - with pytest.raises((ValueError, TypeError), match="object"): + # cupy doesn't support object dtypes. We don't care what the exception + # is here, just that one is raised. + with pytest.raises(Exception, match="object"): check_array(df, mem_type=mem_type) else: # dtype=None does no conversion by default From a4c824be458ea73af2abb4afa2efb21d5c963a17 Mon Sep 17 00:00:00 2001 From: James Lamb Date: Sat, 16 May 2026 07:13:52 -0500 Subject: [PATCH 06/17] loosen threshold in test_mbsgd_regressor test (#8122) Closes #8121 Proposes slightly reducing the threshold in the R-squared check for `MBSGDRegressor` tests. Authors: - James Lamb (https://github.com/jameslamb) Approvers: - Simon Adorf (https://github.com/csadorf) URL: https://github.com/rapidsai/cuml/pull/8122 --- python/cuml/tests/test_mbsgd_regressor.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/cuml/tests/test_mbsgd_regressor.py b/python/cuml/tests/test_mbsgd_regressor.py index 6d3da2b9ef..02a2619605 100644 --- a/python/cuml/tests/test_mbsgd_regressor.py +++ b/python/cuml/tests/test_mbsgd_regressor.py @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 import cupy as cp import numpy as np @@ -138,7 +138,7 @@ def test_mbsgd_regressor(lrate, penalty, make_dataset): cu_pred = model.predict(X_test) cu_r2 = r2_score(cu_pred, y_test) - assert cu_r2 >= 0.88 + assert cu_r2 >= 0.87 def test_mbsgd_regressor_default(make_dataset): From 8e263202f932c572df85ab1b74587556cc96492c Mon Sep 17 00:00:00 2001 From: Tim Head Date: Mon, 18 May 2026 22:10:47 +0200 Subject: [PATCH 07/17] DOC Add hyper-parameter search example (#8095) This adds a cuml.accel example that illustrates the point that using cuml.accel makes it easier to do things because you aren't interrupted by things taking forever and forever. It uses a simple pipeline that can be fully GPU accelerated and then searches a few hyper-parameter combinations for that pipeline. Authors: - Tim Head (https://github.com/betatim) Approvers: - Jim Crist-Harif (https://github.com/jcrist) URL: https://github.com/rapidsai/cuml/pull/8095 --- .../examples/hyperparameter_search.ipynb | 252 ++++++++++++++++++ docs/source/cuml-accel/examples/index.rst | 1 + 2 files changed, 253 insertions(+) create mode 100644 docs/source/cuml-accel/examples/hyperparameter_search.ipynb diff --git a/docs/source/cuml-accel/examples/hyperparameter_search.ipynb b/docs/source/cuml-accel/examples/hyperparameter_search.ipynb new file mode 100644 index 0000000000..93a485db7e --- /dev/null +++ b/docs/source/cuml-accel/examples/hyperparameter_search.ipynb @@ -0,0 +1,252 @@ +{ + "cells": [ + { + "cell_type": "markdown", + "id": "ea791931", + "metadata": {}, + "source": [ + "# Hyperparameter Search with RandomizedSearchCV\n", + "\n", + "This notebook demonstrates how `cuml.accel` speeds up a hyperparameter search\n", + "workflow. Having your train of thought interrupted by long running steps in\n", + "a workflow is not great. By using `cuml.accel` you can take a workflow that\n", + "is tedious because it takes minutes to complete and make it complete in 30s.\n", + "\n", + "In this example we build a preprocessing + classification pipeline and use\n", + "`RandomizedSearchCV` to find the best configuration. However, the principle\n", + "of using `cuml.accel` to take a task from \"requires a coffee break per\n", + "iteration\" to \"it is fun to iterate on ideas\" by speeding it up applies\n", + "to many other tasks as well.\n", + "\n", + "**Pipeline:** `StandardScaler` → `PCA` → `KNeighborsClassifier`\n", + "\n", + "KNN is distance-based, so the preprocessing steps are essential:\n", + "- `StandardScaler` normalises features that span very different ranges\n", + " (elevation 0–3800 vs binary soil-type indicators 0/1).\n", + "- `PCA` reduces the 54-dimensional feature space (40 of which are sparse\n", + " one-hot columns) to a compact representation where distances are more\n", + " informative.\n", + "\n", + "**Dataset:** Forest Cover Type (300K subsample, 54 features, 7 classes).\n", + "\n", + "Without `cuml.accel`, this search takes several minutes (CPU,\n", + "`n_jobs=10`). With `cuml.accel` enabled the same search completes in\n", + "under a minute.\n", + "\n", + "All three pipeline steps (`StandardScaler`, `PCA`, `KNeighborsClassifier`)\n", + "are GPU-accelerated by `cuml.accel`." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "557d6f89", + "metadata": {}, + "outputs": [], + "source": [ + "%load_ext cuml.accel" + ] + }, + { + "cell_type": "markdown", + "id": "18f941fe", + "metadata": {}, + "source": [ + "## Load and prepare the dataset\n", + "\n", + "We use the [Forest Cover Type](https://archive.ics.uci.edu/dataset/31/covertype)\n", + "dataset (581K samples, 54 features, 7 cover-type classes). To keep runtimes\n", + "manageable we subsample to 300K rows and split 80/20 into train and test sets." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "7ae45ccc", + "metadata": {}, + "outputs": [], + "source": [ + "import numpy as np\n", + "from sklearn.datasets import fetch_covtype\n", + "from sklearn.model_selection import train_test_split\n", + "\n", + "X_full, y_full = fetch_covtype(return_X_y=True)\n", + "\n", + "N_SUBSAMPLE = 300_000\n", + "rng = np.random.RandomState(42)\n", + "idx = rng.choice(len(X_full), size=N_SUBSAMPLE, replace=False)\n", + "X, y = X_full[idx], y_full[idx]\n", + "\n", + "X_train, X_test, y_train, y_test = train_test_split(\n", + " X, y, test_size=0.2, random_state=42, stratify=y,\n", + ")\n", + "\n", + "print(f\"Full dataset: {X_full.shape[0]:,} samples, {X_full.shape[1]} features\")\n", + "print(f\"Subsample: {N_SUBSAMPLE:,}\")\n", + "print(f\"Train: {X_train.shape[0]:,}\")\n", + "print(f\"Test: {X_test.shape[0]:,}\")\n", + "print(f\"Classes: {len(np.unique(y_train))}\")" + ] + }, + { + "cell_type": "markdown", + "id": "0f6b0b40", + "metadata": {}, + "source": [ + "## Define the pipeline and search space\n", + "\n", + "The pipeline chains three steps, each GPU-accelerated by `cuml.accel`:\n", + "\n", + "1. `StandardScaler` — normalise feature scales so that distance computations\n", + " are not dominated by high-magnitude features like elevation.\n", + "2. `PCA` — project the 54 features (many of which are sparse one-hot\n", + " indicators) into a lower-dimensional space.\n", + "3. `KNeighborsClassifier` — classify based on nearest neighbours in the\n", + " PCA-reduced space.\n", + "\n", + "We search over PCA dimensionality, number of neighbours, distance weighting,\n", + "and distance metric." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "8fe0d3fd", + "metadata": {}, + "outputs": [], + "source": [ + "from scipy.stats import randint\n", + "from sklearn.decomposition import PCA\n", + "from sklearn.neighbors import KNeighborsClassifier\n", + "from sklearn.pipeline import Pipeline\n", + "from sklearn.preprocessing import StandardScaler\n", + "\n", + "pipe = Pipeline([\n", + " (\"scaler\", StandardScaler()),\n", + " (\"pca\", PCA()),\n", + " (\"knn\", KNeighborsClassifier()),\n", + "])\n", + "\n", + "param_distributions = {\n", + " \"pca__n_components\": [10, 20, 30, 40],\n", + " \"knn__n_neighbors\": randint(3, 30),\n", + " \"knn__weights\": [\"uniform\", \"distance\"],\n", + " \"knn__metric\": [\"euclidean\", \"manhattan\"],\n", + "}" + ] + }, + { + "cell_type": "markdown", + "id": "daf61609", + "metadata": {}, + "source": [ + "## Run the search\n", + "\n", + "We sample 20 random parameter combinations and evaluate each with 5-fold\n", + "cross-validation, for a total of 100 pipeline fits. With `cuml.accel` active\n", + "this takes ~30 seconds; without it (CPU, `n_jobs=10`) the same search takes\n", + "~4.5 minutes." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "1163fa12", + "metadata": {}, + "outputs": [], + "source": [ + "%%time\n", + "\n", + "from sklearn.model_selection import RandomizedSearchCV\n", + "\n", + "search = RandomizedSearchCV(\n", + " pipe,\n", + " param_distributions,\n", + " n_iter=20,\n", + " cv=5,\n", + " scoring=\"accuracy\",\n", + " random_state=42,\n", + " # For CPU, set n_jobs to a higher number\n", + " n_jobs=1,\n", + " refit=True,\n", + ")\n", + "search.fit(X_train, y_train)" + ] + }, + { + "cell_type": "markdown", + "id": "110d5cba", + "metadata": {}, + "source": [ + "## Inspect the results\n", + "\n", + "Let's look at the best hyperparameters found by the search and how the\n", + "top configurations compare." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "ce51ad16", + "metadata": {}, + "outputs": [], + "source": [ + "print(\"Best parameters:\")\n", + "for param, val in sorted(search.best_params_.items()):\n", + " print(f\" {param}: {val}\")\n", + "print(f\"\\nBest CV accuracy: {search.best_score_:.4f}\")" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "02167a16", + "metadata": {}, + "outputs": [], + "source": [ + "import pandas as pd\n", + "\n", + "cv = pd.DataFrame(search.cv_results_)\n", + "cv = cv.sort_values(\"rank_test_score\")\n", + "cv[[\"param_pca__n_components\", \"param_knn__n_neighbors\",\n", + " \"param_knn__weights\", \"param_knn__metric\",\n", + " \"mean_test_score\", \"std_test_score\", \"mean_fit_time\"]].head(10)" + ] + }, + { + "cell_type": "markdown", + "id": "d3f26add", + "metadata": {}, + "source": [ + "## Evaluate on the test set\n", + "\n", + "`RandomizedSearchCV` with `refit=True` automatically refits the best model on\n", + "the full training set. We can use it directly to score on held-out data." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "20befa55", + "metadata": {}, + "outputs": [], + "source": [ + "test_acc = search.score(X_test, y_test)\n", + "print(f\"Test accuracy: {test_acc:.4f}\")" + ] + } + ], + "metadata": { + "kernelspec": { + "display_name": "Python 3", + "language": "python", + "name": "python3" + }, + "language_info": { + "name": "python", + "version": "3.11.0" + } + }, + "nbformat": 4, + "nbformat_minor": 5 +} diff --git a/docs/source/cuml-accel/examples/index.rst b/docs/source/cuml-accel/examples/index.rst index 3ea9b53186..9318614c3f 100644 --- a/docs/source/cuml-accel/examples/index.rst +++ b/docs/source/cuml-accel/examples/index.rst @@ -13,3 +13,4 @@ examples in this section is available in the cuML GitHub repository at `examples plot_kmeans_digits.ipynb third-party-apps.rst onnx_export.ipynb + hyperparameter_search.ipynb From fcf714d159be407655dbcdcd671942684c07fee0 Mon Sep 17 00:00:00 2001 From: Jim Crist-Harif Date: Mon, 18 May 2026 16:25:11 -0500 Subject: [PATCH 08/17] xfail `test_onnx[RandomForestClassifier]` (#8127) This has started to fail. xfailing for now until the issue can be investigated. Stopgap for #8125. Fixes #8129. Authors: - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Simon Adorf (https://github.com/csadorf) URL: https://github.com/rapidsai/cuml/pull/8127 --- python/cuml/cuml_accel_tests/test_onnx.py | 4 ++++ python/cuml/tests/test_incremental_pca.py | 2 +- 2 files changed, 5 insertions(+), 1 deletion(-) diff --git a/python/cuml/cuml_accel_tests/test_onnx.py b/python/cuml/cuml_accel_tests/test_onnx.py index 7f279ad936..d46c320473 100644 --- a/python/cuml/cuml_accel_tests/test_onnx.py +++ b/python/cuml/cuml_accel_tests/test_onnx.py @@ -83,6 +83,10 @@ def regression_data(): pytest.param( RandomForestClassifier(n_estimators=20, max_depth=8, random_state=42), id="RandomForestClassifier", + marks=pytest.mark.xfail( + reason="New failure, see https://github.com/rapidsai/cuml/issues/8125", + strict=False, + ), ), pytest.param(KNeighborsClassifier(), id="KNeighborsClassifier"), pytest.param(LinearSVC(dual="auto"), id="LinearSVC"), diff --git a/python/cuml/tests/test_incremental_pca.py b/python/cuml/tests/test_incremental_pca.py index f831eb8718..021e0dd057 100644 --- a/python/cuml/tests/test_incremental_pca.py +++ b/python/cuml/tests/test_incremental_pca.py @@ -109,7 +109,7 @@ def test_partial_fit( sk_t = sk_ipca.transform(X) sk_inv = sk_ipca.inverse_transform(sk_t) - assert array_equal(cu_inv, sk_inv, 6e-5, with_sign=True) + assert array_equal(cu_inv, sk_inv, 1e-3, with_sign=True) def test_exceptions(): From ee64decb921d5312c57a32fd1671282ad0f86f6a Mon Sep 17 00:00:00 2001 From: Jim Crist-Harif Date: Mon, 18 May 2026 19:30:55 -0500 Subject: [PATCH 09/17] Mark `test_precomputed_sparse_transform_on_iris` as flaky (#8130) The dense version is already marked as flaky, but I've now seen the sparse version fail twice. Authors: - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Simon Adorf (https://github.com/csadorf) URL: https://github.com/rapidsai/cuml/pull/8130 --- python/cuml/cuml_accel_tests/upstream/umap/xfail-list.yaml | 1 + 1 file changed, 1 insertion(+) diff --git a/python/cuml/cuml_accel_tests/upstream/umap/xfail-list.yaml b/python/cuml/cuml_accel_tests/upstream/umap/xfail-list.yaml index 170b8fefb6..742e2a3d15 100644 --- a/python/cuml/cuml_accel_tests/upstream/umap/xfail-list.yaml +++ b/python/cuml/cuml_accel_tests/upstream/umap/xfail-list.yaml @@ -37,6 +37,7 @@ marker: cuml_accel_flaky strict: false tests: + - "umap.tests.test_umap_on_iris::test_precomputed_sparse_transform_on_iris" - "umap.tests.test_umap_on_iris::test_precomputed_transform_on_iris" - "umap.tests.test_umap_ops::test_umap_transform_embedding_stability" - reason: UMAP using removed numpy function `in1d` From 673aa2746e175b50d60f94c8d144aacf42ca171f Mon Sep 17 00:00:00 2001 From: Jim Crist-Harif Date: Mon, 18 May 2026 21:15:06 -0500 Subject: [PATCH 10/17] A few fixes for sklearn 1.9 pre-release (#8126) This addresses failures in `python/cuml/tests` and `python/cuml/cuml_accel_tests` when run with the most recent sklearn 1.9 pre-release. With one small exception, this just required a few tweaks to some tests. I did not address any failures when running the upstream sklearn test suite with `cuml.accel`, as that would require much deeper changes to setup xfails. Authors: - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Simon Adorf (https://github.com/csadorf) URL: https://github.com/rapidsai/cuml/pull/8126 --- .../cuml/accel/_overrides/sklearn/preprocessing.py | 8 ++++---- .../cuml/cuml_accel_tests/integration/test_svc.py | 6 ++++++ .../cuml/cuml_accel_tests/integration/test_tsvd.py | 14 ++++++++++++-- python/cuml/tests/test_sklearn_compatibility.py | 5 ----- python/cuml/tests/test_sklearn_import_export.py | 3 +++ 5 files changed, 25 insertions(+), 11 deletions(-) diff --git a/python/cuml/cuml/accel/_overrides/sklearn/preprocessing.py b/python/cuml/cuml/accel/_overrides/sklearn/preprocessing.py index 24940a9e6f..389b9634dc 100644 --- a/python/cuml/cuml/accel/_overrides/sklearn/preprocessing.py +++ b/python/cuml/cuml/accel/_overrides/sklearn/preprocessing.py @@ -76,13 +76,13 @@ def _check_targetencoder_y(y): class TargetEncoder(ProxyBase): _gpu_class = cuml.preprocessing.TargetEncoder - def _gpu_fit(self, X, y, **kwargs): + def _gpu_fit(self, X, y): _check_targetencoder_y(y) - return self._gpu.fit(X, y, **kwargs) + return self._gpu.fit(X, y) - def _gpu_fit_transform(self, X, y, **kwargs): + def _gpu_fit_transform(self, X, y, **params): _check_targetencoder_y(y) - return self._gpu.fit_transform(X, y, **kwargs) + return self._gpu.fit_transform(X, y, **params) def _gpu_get_feature_names_out(self, input_features=None): """Return feature names for output features. diff --git a/python/cuml/cuml_accel_tests/integration/test_svc.py b/python/cuml/cuml_accel_tests/integration/test_svc.py index ca3de73381..4f7c480803 100644 --- a/python/cuml/cuml_accel_tests/integration/test_svc.py +++ b/python/cuml/cuml_accel_tests/integration/test_svc.py @@ -37,6 +37,12 @@ def test_svc(binary): assert svc.score(X, y) > 0.5 +@pytest.mark.filterwarnings( + "ignore:The `probability` parameter was deprecated:FutureWarning" +) +@pytest.mark.filterwarnings( + "ignore:Attribute `prob[AB]_` was deprecated:FutureWarning" +) def test_svc_probability(binary): X, y = binary svc = SVC(probability=True).fit(X, y) diff --git a/python/cuml/cuml_accel_tests/integration/test_tsvd.py b/python/cuml/cuml_accel_tests/integration/test_tsvd.py index 3912d6c0e3..eb87d193a0 100644 --- a/python/cuml/cuml_accel_tests/integration/test_tsvd.py +++ b/python/cuml/cuml_accel_tests/integration/test_tsvd.py @@ -1,10 +1,12 @@ # -# SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # import numpy as np import pytest +import sklearn +from packaging.version import Version from scipy.sparse import csr_matrix from sklearn.datasets import make_classification from sklearn.decomposition import TruncatedSVD @@ -100,11 +102,19 @@ def test_truncated_svd_tol(svd_data, tol): @pytest.mark.parametrize( - "power_iteration_normalizer", ["auto", "OR", "LU", "none"] + "power_iteration_normalizer", ["auto", "QR", "LU", "none"] ) def test_truncated_svd_power_iteration_normalizer( svd_data, power_iteration_normalizer ): + if ( + Version(sklearn.__version__) < Version("1.9.0.dev0") + and power_iteration_normalizer == "QR" + ): + pytest.skip( + "power_iteration_normalizer 'QR' is not supported in scikit-learn < 1.9.0" + ) + X, _ = svd_data svd = TruncatedSVD( n_components=10, diff --git a/python/cuml/tests/test_sklearn_compatibility.py b/python/cuml/tests/test_sklearn_compatibility.py index 58d8f50d7b..17a96f8c5d 100644 --- a/python/cuml/tests/test_sklearn_compatibility.py +++ b/python/cuml/tests/test_sklearn_compatibility.py @@ -208,7 +208,6 @@ }, KernelDensity: { "check_estimator_tags_renamed": "No support for modern tags infrastructure", - "check_all_zero_sample_weights_error": "KernelDensity does not validate all-zero sample weights", }, EmpiricalCovariance: { "check_estimator_tags_renamed": "No support for modern tags infrastructure", @@ -247,7 +246,6 @@ "check_sample_weights_not_an_array": "sample_weight not implemented", "check_sample_weights_shape": "sample_weight not implemented", "check_sample_weight_equivalence_on_dense_data": "sample_weight not implemented", - "check_all_zero_sample_weights_error": "sample_weight not implemented", "check_sample_weights_list": "sample_weight not implemented", "check_sample_weights_not_overwritten": "sample_weight not implemented", "check_sample_weight_equivalence_on_sparse_data": "sample_weight not implemented", @@ -260,7 +258,6 @@ "check_sample_weights_not_an_array": "sample_weight not implemented", "check_sample_weights_shape": "sample_weight not implemented", "check_sample_weight_equivalence_on_dense_data": "sample_weight not implemented", - "check_all_zero_sample_weights_error": "sample_weight not implemented", "check_sample_weights_list": "sample_weight not implemented", "check_sample_weights_not_overwritten": "sample_weight not implemented", "check_sample_weight_equivalence_on_sparse_data": "sample_weight not implemented", @@ -273,7 +270,6 @@ "check_sample_weights_not_an_array": "sample_weight not implemented", "check_sample_weights_shape": "sample_weight not implemented", "check_sample_weight_equivalence_on_dense_data": "sample_weight not implemented", - "check_all_zero_sample_weights_error": "sample_weight not implemented", "check_sample_weights_list": "sample_weight not implemented", "check_sample_weights_not_overwritten": "sample_weight not implemented", "check_sample_weight_equivalence_on_sparse_data": "sample_weight not implemented", @@ -286,7 +282,6 @@ "check_sample_weights_not_an_array": "sample_weight not implemented", "check_sample_weights_shape": "sample_weight not implemented", "check_sample_weight_equivalence_on_dense_data": "sample_weight not implemented", - "check_all_zero_sample_weights_error": "sample_weight not implemented", "check_sample_weights_list": "sample_weight not implemented", "check_sample_weights_not_overwritten": "sample_weight not implemented", "check_sample_weight_equivalence_on_sparse_data": "sample_weight not implemented", diff --git a/python/cuml/tests/test_sklearn_import_export.py b/python/cuml/tests/test_sklearn_import_export.py index e3ca75b356..c9a0d09319 100644 --- a/python/cuml/tests/test_sklearn_import_export.py +++ b/python/cuml/tests/test_sklearn_import_export.py @@ -379,6 +379,9 @@ def test_svr(random_state, sparse, kernel): @pytest.mark.filterwarnings( "ignore:The `probability` parameter was deprecated:FutureWarning" ) +@pytest.mark.filterwarnings( + "ignore:Attribute `prob[AB]_` was deprecated:FutureWarning" +) @pytest.mark.parametrize("sparse", [False, True]) @pytest.mark.parametrize("probability", [False, True]) @pytest.mark.parametrize("kernel", ["rbf", "precomputed"]) From ac619316844f1ecd4b00ca74f8d86dcb6dd23845 Mon Sep 17 00:00:00 2001 From: Jim Crist-Harif Date: Mon, 18 May 2026 21:15:29 -0500 Subject: [PATCH 11/17] Cleanup LabelBinarizer (#8101) - Updates `LabelBinarizer` to follow standard cuml and sklearn conventions (simple `__init__`, no mutation, type reflection, ...) - Applies new validation - Adds `sparse_input_` and `y_type_` attributes - Improves validation and error messages - Improves resilience and sklearn compatibility - Improves test coverage - Adds support for sklearn interop - Adds support for cuml.accel - Improves docstrings This required one change to `cuml.internals.validation` around handling of unsupported dtypes for `cupyx.scipy.sparse`. This is split out into a separate commit with a new test case. Part of #7317. Fixes #8087. Authors: - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Simon Adorf (https://github.com/csadorf) URL: https://github.com/rapidsai/cuml/pull/8101 --- docs/source/cuml-accel/faq.rst | 1 + docs/source/cuml-accel/limitations.rst | 5 + .../accel/_overrides/sklearn/preprocessing.py | 8 + python/cuml/cuml/internals/validation.py | 50 +- python/cuml/cuml/preprocessing/label.py | 605 +++++++++++++----- .../integration/test_preprocessing.py | 40 ++ python/cuml/tests/test_label_binarizer.py | 274 +++++++- python/cuml/tests/test_naive_bayes.py | 19 +- .../cuml/tests/test_sklearn_import_export.py | 24 + python/cuml/tests/test_validation.py | 58 +- 10 files changed, 860 insertions(+), 224 deletions(-) diff --git a/docs/source/cuml-accel/faq.rst b/docs/source/cuml-accel/faq.rst index 87ccf0591a..e83b724749 100644 --- a/docs/source/cuml-accel/faq.rst +++ b/docs/source/cuml-accel/faq.rst @@ -75,6 +75,7 @@ the following estimators are mostly or entirely accelerated when run with * ``sklearn.preprocessing.MaxAbsScaler`` * ``sklearn.preprocessing.PolynomialFeatures`` * ``sklearn.preprocessing.LabelEncoder`` + * ``sklearn.preprocessing.LabelBinarizer`` * ``sklearn.preprocessing.TargetEncoder`` * ``sklearn.svm.SVC`` * ``sklearn.svm.SVR`` diff --git a/docs/source/cuml-accel/limitations.rst b/docs/source/cuml-accel/limitations.rst index fa0b991c9f..5fa70dc860 100644 --- a/docs/source/cuml-accel/limitations.rst +++ b/docs/source/cuml-accel/limitations.rst @@ -463,6 +463,11 @@ LabelEncoder ``LabelEncoder`` supports all cases and will never fall back to CPU. +LabelBinarizer +^^^^^^^^^^^^^^ + +``LabelBinarizer`` supports all cases and will never fall back to CPU. + TargetEncoder ^^^^^^^^^^^^^ diff --git a/python/cuml/cuml/accel/_overrides/sklearn/preprocessing.py b/python/cuml/cuml/accel/_overrides/sklearn/preprocessing.py index 389b9634dc..cd659776d5 100644 --- a/python/cuml/cuml/accel/_overrides/sklearn/preprocessing.py +++ b/python/cuml/cuml/accel/_overrides/sklearn/preprocessing.py @@ -19,6 +19,7 @@ "PolynomialFeatures", "TargetEncoder", "LabelEncoder", + "LabelBinarizer", ) @@ -58,6 +59,13 @@ class LabelEncoder(ProxyBase): _gpu_class = cuml.preprocessing.LabelEncoder +class LabelBinarizer(ProxyBase): + _gpu_class = cuml.preprocessing.LabelBinarizer + + def _gpu_inverse_transform(self, Y, threshold=None): + return self._gpu.inverse_transform(Y, threshold=threshold) + + def _check_targetencoder_y(y): """Check if inputs are supported on GPU. diff --git a/python/cuml/cuml/internals/validation.py b/python/cuml/cuml/internals/validation.py index e8afe0aa55..4c8d5e1640 100644 --- a/python/cuml/cuml/internals/validation.py +++ b/python/cuml/cuml/internals/validation.py @@ -591,11 +591,31 @@ def check_array( raise ValueError(f"Unsupported {mem_type=!r}") if order not in ("F", "C", "A", None): raise ValueError(f"Unsupported {order=!r}") + if dtype is not None: if not isinstance(dtype, (list, tuple)): dtype = [dtype] dtype = [np.dtype(i) for i in dtype] + is_sparse = cp_sp.issparse(array) or sp.issparse(array) + if is_sparse and ( + mem_type == "device" or (mem_type is None and cp_sp.issparse(array)) + ): + # XXX: cupyx.scipy.sparse doesn't support integral dtypes. If a dtype + # is specified, we filter to only supported types (erroring if no + # supported types specified). If dtype=None, we use the input dtype if + # supported, and the closest floating type otherwise. + if dtype is not None: + if not any(d.kind in "fb" for d in dtype): + raise ValueError( + f"No dtype in {dtype} is supported by cupyx.scipy.sparse" + ) + dtype = [d for d in dtype if d.kind in "fb"] + elif array.dtype.kind not in "fb": + dtype = [ + np.dtype("f4") if array.dtype.itemsize <= 4 else np.dtype("f8") + ] + # Extract original array type and dtype (when possible) array_type = type(array) if isinstance(array, (cudf.DataFrame, pd.DataFrame)): @@ -629,13 +649,13 @@ def check_array( else: dtype = array_dtype elif dtype is not None: - # No original dtype, use first dtype in list inputs + # No original dtype, use first provided dtype dtype = dtype[0] # Coerce `array` to numpy/cupy/scipy.sparse/cupyx.scipy.sparse values as # requested. For dataframe-like inputs also extract the index for later use. index = None - if cp_sp.issparse(array) or sp.issparse(array): + if is_sparse: orig_sparse_array = array # Handle sparse inputs if isinstance(accept_sparse, str): @@ -667,17 +687,23 @@ def check_array( ensure_min_features=ensure_min_features, ) - # Coerce data to accepted dtype if needed - if dtype is not None and array.dtype != dtype: - array = array.astype(dtype) - - # Coerce to device or host if needed - if mem_type == "device" and not cp_sp.issparse(array): - if array.ndim != 2: - raise ValueError("cupyx.scipy.sparse only supports 2D arrays") - array = getattr(cp_sp, f"{array.format}_matrix")(array) - elif mem_type == "host" and not sp.issparse(array): + # Coerce to proper dtype and mem_type if needed + if mem_type == "host" and not sp.issparse(array): + # Coerce to device, then coerce dtype. We do this to save device + # memory, and since scipy supports more dtypes. array = array.get() + if dtype is not None and array.dtype != dtype: + array = array.astype(dtype) + else: + # Otherwise coerce dtype, then mem_type if needed + if dtype is not None and array.dtype != dtype: + array = array.astype(dtype) + if mem_type == "device" and not cp_sp.issparse(array): + if array.ndim != 2: + raise ValueError( + "cupyx.scipy.sparse only supports 2D arrays" + ) + array = getattr(cp_sp, f"{array.format}_matrix")(array) # Copy if needed if copy and array is orig_sparse_array: diff --git a/python/cuml/cuml/preprocessing/label.py b/python/cuml/cuml/preprocessing/label.py index a042130c8e..ffaea9c032 100644 --- a/python/cuml/cuml/preprocessing/label.py +++ b/python/cuml/cuml/preprocessing/label.py @@ -1,82 +1,272 @@ # SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # - +import cudf import cupy as cp -import cupyx -import scipy.sparse +import cupyx.scipy.sparse as cp_sp +import numpy as np +import scipy.sparse as sp import cuml.internals -from cuml.common import CumlArray -from cuml.common.array_descriptor import CumlArrayDescriptor -from cuml.internals.array_sparse import SparseCumlArray +from cuml.common.classification import decode_labels +from cuml.internals.array import CumlArray from cuml.internals.base import Base -from cuml.prims.label import make_monotonic +from cuml.internals.interop import InteropMixin +from cuml.internals.validation import ( + check_array, + check_classification_targets, + check_is_fitted, +) + + +def _label_binarize( + y, + *, + classes=..., + neg_label=0, + pos_label=1, + sparse_output=False, + accept_multilabel=True, +): + """A helper used to implement `label_binarize` and `LabelBinarizer`""" + if neg_label >= pos_label: + raise ValueError( + f"{neg_label=} must be strictly less than {pos_label=}." + ) + if sparse_output and (pos_label == 0 or neg_label != 0): + raise ValueError( + "Sparse binarization is only supported with non " + "zero pos_label and zero neg_label, got " + f"{pos_label=} and {neg_label=}" + ) -@cuml.internals.reflect -def label_binarize( - y, classes, neg_label=0, pos_label=1, sparse_output=False -) -> SparseCumlArray: - """ - A stateless helper function to dummy encode multi-class labels. + if classes is not ...: + if hasattr(classes, "__cuda_array_interface__"): + classes = cp.asarray(classes) + else: + classes = np.asarray(classes) - Parameters - ---------- + # To account for pos_label == 0 in the dense case + if pos_switch := pos_label == 0: + pos_label = -neg_label - y : array-like of size [n_samples,] or [n_samples, n_classes] - classes : the set of unique classes in the input - neg_label : integer the negative value for transformed output - pos_label : integer the positive value for transformed output - sparse_output : bool whether to return sparse array - """ + is_multioutput = False + + if sparse_input := (cp_sp.issparse(y) or sp.issparse(y)): + # Coerce to cupyx.scipy.sparse.csr_matrix + y = check_array( + y, + dtype=("float32", "float64"), + accept_sparse="csr", + ensure_min_samples=0, + ensure_all_finite=False, + ) + # Ensure y is integral and finite + check_classification_targets(y.data) + is_multioutput = len(cp.unique(y.data)) > 2 + else: + # cudf may coerce the dtype, store the original so we can cast back later + input_dtype = y.dtype if isinstance(y, np.ndarray) else None + + if not isinstance(y, (cudf.DataFrame, cudf.Series)): + y = check_array( + y, + mem_type=None, + ensure_2d=False, + ensure_min_samples=0, + ensure_all_finite=False, + input_name="y", + ) + # If no original dtype found on input, use the coerced one instead + if input_dtype is None: + input_dtype = y.dtype - classes = cp.asarray(classes, dtype=classes.dtype) - labels = cp.asarray(y, dtype=y.dtype) + y = (cudf.DataFrame if y.ndim == 2 else cudf.Series)( + y, dtype=(np.dtype("O") if y.dtype.kind in "U" else None) + ) + else: + y = y.reset_index(drop=True) + + if y.ndim == 2 and y.shape[1] == 1: + y = y.iloc[:, 0] + if y.ndim == 1: + check_classification_targets(y) + elif y.ndim == 2: + if y.select_dtypes(exclude=["number", "bool"]).shape[1]: + is_multioutput = True + else: + y = y.to_cupy() + check_classification_targets(y) + is_multioutput = len(cp.unique(y)) > 2 + + if is_multioutput: + raise ValueError( + "Multioutput target data is not supported with label binarization" + ) - # Check that all labels are in classes - if not bool(cp.all(cp.isin(labels, classes))): - raise ValueError("Unseen classes encountered in input") + if y.shape[0] == 0: + raise ValueError("y has 0 samples, while a minimum of 1 is required") + + has_unseen = False + if y.ndim == 1: + # binary or multiclass + # y is a cudf.Series + y = y.astype("category") + if classes is ...: + classes = y.cat.categories + # XXX: cudf's to_numpy doesn't support conversions for all + # dtypes. Roundtrip through object dtype when necessary. + try: + classes = classes.to_numpy(dtype=input_dtype) + except NotImplementedError: + classes = classes.to_numpy(dtype="object") + # cudf will sometimes translate non-numeric dtypes. Coerce back to + # the input dtype if the input was originally a numpy array. + if input_dtype is not None: + classes = classes.astype(input_dtype, copy=False) + indices = cp.asarray(y.cat.codes) + indptr = cp.arange(len(y) + 1) + else: + y = y.cat.set_categories(classes) + if has_unseen := y.has_nulls: + mask = ~y.isnull() + indices = cp.asarray(y[mask].cat.codes) + indptr = cp.concatenate([cp.array([0]), mask.cumsum()]) + else: + indices = cp.asarray(y.cat.codes) + indptr = cp.arange(len(y) + 1) + + if len(classes) == 1: + # Special case binary with 1 class -> neg_label + y_type = "binary" + out = cp_sp.csr_matrix((len(y), 1), dtype="float32") + else: + y_type = "binary" if len(classes) <= 2 else "multiclass" + data = cp.full(len(indices), pos_label, dtype="float32") + out = cp_sp.csr_matrix( + (data, indices, indptr), shape=(len(y), len(classes)) + ) + if not sparse_output: + out = out.toarray() + else: + # multilabel-indicator + # y is a cupy.ndarray or cupyx.scipy.sparse.csr_matrix + y_type = "multilabel-indicator" - row_ind = cp.arange(0, labels.shape[0], 1, dtype=y.dtype) - col_ind, _ = make_monotonic(labels, classes, copy=True) + if not accept_multilabel: + raise ValueError( + "The object was not fitted with multilabel input." + ) - val = cp.full(row_ind.shape[0], pos_label, dtype=y.dtype) + if classes is ...: + classes = np.arange(y.shape[1]) + elif len(classes) != y.shape[1]: + raise ValueError( + f"classes {classes} mismatch with the labels " + f"{np.arange(y.shape[1])} found in the data" + ) - sp = cupyx.scipy.sparse.coo_matrix( - (val, (row_ind, col_ind)), - shape=(col_ind.shape[0], classes.shape[0]), - dtype=cp.float32, - ) + if sparse_output: + out = cp_sp.csr_matrix(y.astype("float32")) + if pos_label != 1: + out.data = cp.full_like(out.data, pos_label) + else: + out = y.toarray() if cp_sp.issparse(y) else y.copy() + if pos_label != 1: + out[out != 0] = pos_label - is_binary = classes.shape[0] == 2 + if not sparse_output: + if neg_label != 0: + out[out == 0] = neg_label - if sparse_output: - sp = sp.tocsr() - if is_binary: - sp = sp.getcol(1) # getcol does not support -1 indexing - return sp - else: - arr = sp.toarray().astype(y.dtype) - arr[arr == 0] = neg_label - if is_binary: - arr = arr[:, -1].reshape((-1, 1)) - return arr + if pos_switch: + out[out == pos_label] = 0 + + out = out.astype("int32", copy=False) + + # XXX: In a binary problem with unseen labels we return a matrix of shape + # (n_samples, 2), while sklearn returns (n_samples, 1). This is an edge + # case (binary inputs for `LabelBinarizer` are a bit odd, as are unseen + # classes. We view the sklearn behavior as a bug (see + # https://github.com/scikit-learn/scikit-learn/issues/13674), since with + # their encoding unseen labels are conflated with label 0 rather than + # encoded as missing via all 0s (as in the multiclass case). + if y_type == "binary" and not has_unseen: + out = out[:, [-1]] + return out, classes, y_type, sparse_input -class LabelBinarizer(Base): + +@cuml.internals.reflect +def label_binarize(y, classes, neg_label=0, pos_label=1, sparse_output=False): """ - A multi-class dummy encoder for labels. + Binarize labels in a one-vs-all fashion. Parameters ---------- + y : array-like or sparse matrix, shape (n_samples,) or (n_samples, n_classes) + Target values. The 2-d matrix should only contain 0 and 1, in the + multilabel-indicator format. + classes : array-like of shape (n_classes,) + The class labels for each class. + neg_label : int, default=0 + The value to use for encoding negative labels. + pos_label : int, default=1 + The value to use for encoding positive labels. + sparse_output : bool, default=False + If true, a sparse CSR matrix is returned. + + Returns + ------- + y : array or sparse matrix, shape (n_samples, n_classes) + The encoded labels. Will be a sparse matrix if ``sparse_output=True``. + Shape will be (n_samples, n_classes) for multiclass problems, + (n_samples, 1) for binary problems with no unseen classes, and + (n_samples, 2) for binary problems with unseen classes (a minor, + intentional deviation from sklearn). + + See Also + -------- + LabelBinarizer : A class version of this function. + + Examples + -------- + >>> from cuml.preprocessing import label_binarize + >>> label_binarize([1, 6], classes=[1, 2, 4, 6]) + array([[1, 0, 0, 0], + [0, 0, 0, 1]], dtype=int32) + + Binary targets result in a column vector: + + >>> label_binarize(['a', 'b', 'b', 'a'], classes=['a', 'b']) + array([[0], + [1], + [1], + [0]], dtype=int32) + """ + out, _, _, _ = _label_binarize( + y, + classes=classes, + neg_label=neg_label, + pos_label=pos_label, + sparse_output=sparse_output, + ) + return out + - neg_label : integer (default=0) - label to be used as the negative binary label - pos_label : integer (default=1) - label to be used as the positive binary label - sparse_output : bool (default=False) - whether to return sparse arrays for transformed output +class LabelBinarizer(Base, InteropMixin): + """ + Binarize labels in a one-vs-all fashion. + + Parameters + ---------- + neg_label : int, default=0 + The value to use for encoding negative labels. + pos_label : int, default=1 + The value to use for encoding positive labels. + sparse_output : bool, default=False + If true, a sparse CSR matrix is returned from ``transform``. verbose : int or boolean, default=False Sets logging level. It must be one of `cuml.common.logger.level_*`. See :ref:`verbosity-levels` for more info. @@ -87,43 +277,43 @@ class LabelBinarizer(Base): (`cuml.global_settings.output_type`) will be used. See :ref:`output-data-type-configuration` for more info. - Examples + Attributes + ---------- + classes_ : numpy.ndarray of shape (n_classes,) + Holds the label for each class. + y_type_ : {'binary', 'multiclass', 'multilabel-indicator'} + The type of the target data. + sparse_input_ : bool + Whether the input data to `fit` was a sparse matrix. + + See Also -------- + label_binarize : A function version of this class. - Create an array with labels and dummy encode them - - .. code-block:: python - - >>> import cupy as cp - >>> import cupyx - >>> from cuml.preprocessing import LabelBinarizer - - >>> labels = cp.asarray([0, 5, 10, 7, 2, 4, 1, 0, 0, 4, 3, 2, 1], - ... dtype=cp.int32) - - >>> lb = LabelBinarizer() - >>> encoded = lb.fit_transform(labels) - >>> print(str(encoded)) - [[1 0 0 0 0 0 0 0] - [0 0 0 0 0 1 0 0] - [0 0 0 0 0 0 0 1] - [0 0 0 0 0 0 1 0] - [0 0 1 0 0 0 0 0] - [0 0 0 0 1 0 0 0] - [0 1 0 0 0 0 0 0] - [1 0 0 0 0 0 0 0] - [1 0 0 0 0 0 0 0] - [0 0 0 0 1 0 0 0] - [0 0 0 1 0 0 0 0] - [0 0 1 0 0 0 0 0] - [0 1 0 0 0 0 0 0]] - >>> decoded = lb.inverse_transform(encoded) - >>> print(str(decoded)) - [ 0 5 10 7 2 4 1 0 0 4 3 2 1] - + Examples + -------- + >>> import cupy as cp + >>> from cuml.preprocessing import LabelBinarizer + >>> y = cp.array([1, 2, 6, 4, 2]) + >>> lb = LabelBinarizer().fit(y) + >>> lb.classes_ + array([1, 2, 4, 6]) + >>> lb.transform(cp.array([1, 6])) + array([[1, 0, 0, 0], + [0, 0, 0, 1]], dtype=int32) + + Binary targets result in a column vector: + + >>> import numpy as np + >>> lb = LabelBinarizer() + >>> lb.fit_transform(np.array(['a', 'b', 'b', 'a'])) + array([[0], + [1], + [1], + [0]], dtype=int32) """ - classes_ = CumlArrayDescriptor() + _cpu_class_path = "sklearn.preprocessing.LabelBinarizer" def __init__( self, @@ -135,128 +325,221 @@ def __init__( output_type=None, ): super().__init__(verbose=verbose, output_type=output_type) - - if neg_label >= pos_label: - raise ValueError( - "neg_label=%s must be less " - "than pos_label=%s." % (neg_label, pos_label) - ) - - if sparse_output and (pos_label == 0 or neg_label != 0): - raise ValueError( - "Sparse binarization is only supported " - "with non-zero" - "pos_label and zero neg_label, got pos_label=%s " - "and neg_label=%s" % (pos_label, neg_label) - ) - self.neg_label = neg_label self.pos_label = pos_label self.sparse_output = sparse_output - self.classes_ = None - @cuml.internals.reflect(reset="type") + @classmethod + def _get_param_names(cls): + return [ + *super()._get_param_names(), + "neg_label", + "pos_label", + "sparse_output", + ] + + def __sklearn_is_fitted__(self) -> bool: + return hasattr(self, "classes_") + + @staticmethod + def _more_static_tags(): + return {"X_types": ["1dlabels"]} + + @classmethod + def _params_from_cpu(cls, model): + return { + "neg_label": model.neg_label, + "pos_label": model.pos_label, + "sparse_output": model.sparse_output, + } + + def _params_to_cpu(self): + return { + "neg_label": self.neg_label, + "pos_label": self.pos_label, + "sparse_output": self.sparse_output, + } + + def _attrs_from_cpu(self, model): + return { + "y_type_": model.y_type_, + "sparse_input_": model.sparse_input_, + "classes_": model.classes_, + } + + def _attrs_to_cpu(self, model): + return { + "y_type_": self.y_type_, + "sparse_input_": self.sparse_input_, + "classes_": self.classes_, + } + + @cuml.internals.run_in_internal_context def fit(self, y) -> "LabelBinarizer": """ - Fit label binarizer + Fit label binarizer. Parameters ---------- y : array of shape [n_samples,] or [n_samples, n_classes] Target values. The 2-d matrix should only contain 0 and 1, - represents multilabel classification. + in the multilabel-indicator format. Returns ------- - self : returns an instance of self. + self : LabelBinarizer + Returns the instance itself. """ - - if y.ndim > 2: - raise ValueError("labels cannot be greater than 2 dimensions") - - if y.ndim == 2: - unique_classes = cp.unique(y) - if unique_classes != [0, 1]: - raise ValueError("2-d array can must be binary") - - self.classes_ = cp.arange(0, y.shape[1]) - else: - self.classes_ = cp.unique(y).astype(y.dtype) - + self.fit_transform(y) return self - @cuml.internals.reflect - def fit_transform(self, y) -> SparseCumlArray: + @cuml.internals.reflect(reset="type") + def fit_transform(self, y): """ - Fit label binarizer and transform multi-class labels to their - dummy-encoded representation. + Fit label binarizer and transform labels to binary labels. Parameters ---------- - y : array of shape [n_samples,] or [n_samples, n_classes] + y : array-like or sparse matrix, shape (n_samples,) or (n_samples, n_classes) + Target values. The 2-d matrix should only contain 0 and 1, in the + multilabel-indicator format. Returns ------- - - arr : array with encoded labels + y : array or sparse matrix + The encoded labels. Shape will be (n_samples, 1) for binary + classification problems. Will be a sparse matrix if + ``sparse_output=True``. """ - return self.fit(y).transform(y) + out, classes, y_type, sparse_input = _label_binarize( + y, + neg_label=self.neg_label, + pos_label=self.pos_label, + sparse_output=self.sparse_output, + ) + self.classes_ = classes + self.y_type_ = y_type + self.sparse_input_ = sparse_input + return out @cuml.internals.reflect - def transform(self, y) -> SparseCumlArray: + def transform(self, y): """ - Transform multi-class labels to their dummy-encoded representation - labels. + Transform labels to binary labels. Parameters ---------- - y : array of shape [n_samples,] or [n_samples, n_classes] + y : array-like or sparse matrix, shape (n_samples,) or (n_samples, n_classes) + Target values. The 2-d matrix should only contain 0 and 1, in the + multilabel-indicator format. Returns ------- - arr : array with encoded labels + y : array or sparse matrix + The encoded labels. Will be a sparse matrix if + ``sparse_output=True``. Shape will be (n_samples, n_classes) for + multiclass problems, (n_samples, 1) for binary problems with no + unseen classes, and (n_samples, 2) for binary problems with unseen + classes (a minor, intentional deviation from sklearn). """ - return label_binarize( + check_is_fitted(self) + out, _, _, _ = _label_binarize( y, - self.classes_, - pos_label=self.pos_label, + classes=self.classes_, neg_label=self.neg_label, + pos_label=self.pos_label, sparse_output=self.sparse_output, + accept_multilabel=self.y_type_ == "multilabel-indicator", ) + return out - @cuml.internals.reflect - def inverse_transform(self, y, *, threshold=None) -> CumlArray: + @cuml.internals.run_in_internal_context + def inverse_transform(self, y, *, threshold=None): """ - Transform binary labels back to original multi-class labels + Transform binary labels back to original labels. Parameters ---------- - - y : array of shape [n_samples, n_classes] - threshold : float this value is currently ignored + y : array-like or sparse matrix, shape (n_samples, n_classes) + The encoded target values. + threshold : float, default=None + Threshold used in the binary and multilabel-indicator cases. + If None, the threshold is assumed to be half way between + ``neg_label`` and ``pos_label``. Returns ------- - - arr : array with original labels + y : array or sparse matrix, shape (n_samples,) or (n_samples, n_classes) + The original target values. """ - # If we are already given multi-class, just return it. - if cupyx.scipy.sparse.isspmatrix(y): - y_mapped = y.tocsr().indices.astype(self.classes_.dtype) - elif scipy.sparse.isspmatrix(y): - y = y.tocsr() - y_mapped = cp.array(y.indices, dtype=y.indices.dtype) - else: - y_mapped = cp.argmax(cp.asarray(y, dtype=y.dtype), axis=1).astype( - y.dtype + check_is_fitted(self) + + # Determine output type + with cuml.internals.exit_internal_context(): + output_type = self._get_output_type(y) + + # Validate and normalize y to a cupy array or csr_matrix + y, index = check_array( + y, + ensure_2d=False, + ensure_min_samples=0, + accept_sparse="csr", + return_index=True, + ) + if y.ndim == 1: + y = y[:, None] + + # Ensure shape is valid with fit classes + n_classes = len(self.classes_) + if n_classes == 2 and y.shape[1] not in (1, 2): + raise ValueError( + f"Expected `y` with 1 or 2 columns, but got {y.shape[1]}" + ) + elif n_classes != 2 and y.shape[1] != n_classes: + raise ValueError( + f"Expected `y` with {n_classes} columns, but got {y.shape[1]}" ) - return CumlArray(data=self.classes_[y_mapped]) + # Transform back to original input + if self.y_type_ == "multiclass": + if cp_sp.issparse(y): + indices = y.argmax(1).flatten() + else: + indices = y.argmax(axis=1) + else: + if threshold is None: + threshold = (self.pos_label + self.neg_label) / 2.0 + + if cp_sp.issparse(y): + if threshold > 0: + indices = (y > threshold).toarray().view("int8") + else: + # Results in a fully dense output, pre-densify to save memory + indices = (y.toarray() > threshold).view("int8") + else: + indices = (y > threshold).view("int8") + + if self.y_type_ == "binary": + if cp_sp.issparse(indices): + indices = indices.toarray() + if y.ndim == 2 and y.shape[1] == 2: + indices = indices[:, 1].flatten() + elif n_classes == 1: + indices = cp.zeros(len(y), dtype="int8") + else: + indices = indices.flatten() + + if self.y_type_ in ("binary", "multiclass"): + return decode_labels( + indices, self.classes_, output_type=output_type, index=index + ) - @classmethod - def _get_param_names(cls): - return super()._get_param_names() + [ - "neg_label", - "pos_label", - "sparse_output", - ] + # For multilabel-indicator we need to handle the conversion manually + if self.sparse_input_: + out = cp_sp.csr_matrix(indices.astype("float32", copy=False)) + if output_type in ("numpy", "pandas"): + out = out.get() + return out + else: + out = indices.astype("int32", copy=False) + return CumlArray(out, index=index).to_output(output_type) diff --git a/python/cuml/cuml_accel_tests/integration/test_preprocessing.py b/python/cuml/cuml_accel_tests/integration/test_preprocessing.py index e2a7a8ff1c..7f2a79464e 100644 --- a/python/cuml/cuml_accel_tests/integration/test_preprocessing.py +++ b/python/cuml/cuml_accel_tests/integration/test_preprocessing.py @@ -4,8 +4,10 @@ import numpy as np import pandas as pd import pytest +import scipy.sparse as sp from sklearn.datasets import make_blobs from sklearn.preprocessing import ( + LabelBinarizer, LabelEncoder, MaxAbsScaler, MinMaxScaler, @@ -118,3 +120,41 @@ def test_label_encoder(): np.testing.assert_array_equal(enc.classes_, np.array(["a", "b"])) y3 = enc.inverse_transform(y2) np.testing.assert_array_equal(y3, y) + + +def test_label_binarizer(): + y = np.array(["a", "b", "a", "c"]) + enc = LabelBinarizer() + + y2 = enc.fit_transform(y) + sol = np.array([[1, 0, 0], [0, 1, 0], [1, 0, 0], [0, 0, 1]]) + np.testing.assert_array_equal(y2, sol) + + np.testing.assert_array_equal(enc.classes_, np.array(["a", "b", "c"])) + assert enc.classes_.dtype == y.dtype + assert enc.y_type_ == "multiclass" + + y3 = enc.transform(np.array(["a", "d"])) + sol = np.array([[1, 0, 0], [0, 0, 0]]) + np.testing.assert_array_equal(y3, sol) + + y4 = enc.inverse_transform(y2) + np.testing.assert_array_equal(y4, y) + + +@pytest.mark.parametrize("sparse", [True, False]) +def test_label_binarizer_multilabel_indicator(sparse): + y = np.array([[1, 0, 0], [0, 1, 0], [0, 0, 1]]) + if sparse: + y = sp.csr_matrix(y) + enc = LabelBinarizer().fit(y) + np.testing.assert_array_equal(enc.classes_, np.array([0, 1, 2])) + assert enc.y_type_ == "multilabel-indicator" + + y2 = enc.inverse_transform(y) + if sparse: + assert isinstance(y2, sp.csr_matrix) + np.testing.assert_array_equal(y.toarray(), y2.toarray()) + else: + assert isinstance(y2, np.ndarray) + np.testing.assert_array_equal(y, y2) diff --git a/python/cuml/tests/test_label_binarizer.py b/python/cuml/tests/test_label_binarizer.py index 78f745823f..d01e606aeb 100644 --- a/python/cuml/tests/test_label_binarizer.py +++ b/python/cuml/tests/test_label_binarizer.py @@ -4,11 +4,9 @@ import cupy as cp import numpy as np import pytest -import scipy.sparse -from sklearn.preprocessing import LabelBinarizer as skLB +import scipy.sparse as sp -from cuml.preprocessing import LabelBinarizer -from cuml.testing.utils import array_equal +from cuml.preprocessing import LabelBinarizer, label_binarize def test_label_binarizer_no_features(): @@ -18,44 +16,262 @@ def test_label_binarizer_no_features(): assert not hasattr(model, "n_features_in_") +def test_label_binarizer_not_fitted(): + lb = LabelBinarizer() + err_msg = "This LabelBinarizer instance is not fitted yet" + with pytest.raises(ValueError, match=err_msg): + lb.transform([]) + with pytest.raises(ValueError, match=err_msg): + lb.inverse_transform([]) + + +def test_label_binarizer_invalid_parameters(): + input_labels = [0, 1, 0, 1] + err_msg = "neg_label=2 must be strictly less than pos_label=1." + lb = LabelBinarizer(neg_label=2, pos_label=1) + with pytest.raises(ValueError, match=err_msg): + lb.fit(input_labels) + err_msg = "neg_label=2 must be strictly less than pos_label=2." + lb = LabelBinarizer(neg_label=2, pos_label=2) + with pytest.raises(ValueError, match=err_msg): + lb.fit(input_labels) + err_msg = ( + "Sparse binarization is only supported with non zero pos_label and zero " + "neg_label, got pos_label=2 and neg_label=1" + ) + lb = LabelBinarizer(neg_label=1, pos_label=2, sparse_output=True) + with pytest.raises(ValueError, match=err_msg): + lb.fit(input_labels) + + +def test_label_binarizer_invalid_y_types(): + err_msg = ( + "Multioutput target data is not supported with label binarization" + ) + with pytest.raises(ValueError, match=err_msg): + LabelBinarizer().fit(np.array([[1, 3], [2, 1]])) + with pytest.raises(ValueError, match=err_msg): + label_binarize(np.array([[1, 3], [2, 1]]), classes=[1, 2, 3]) + + with pytest.raises(ValueError, match="Unknown label type: continuous"): + LabelBinarizer().fit([1.2, 2.7]) + + +def test_label_binarize_mismatch_labels(): + with pytest.raises(ValueError, match="mismatch with the labels"): + label_binarize([[1, 0]], classes=[0, 1, 2]) + + lb = LabelBinarizer().fit(np.array([[0, 0, 1]])) + with pytest.raises(ValueError, match="mismatch with the labels"): + lb.transform(np.array([[0, 1]])) + + +def toarray(x, expect_sparse): + if expect_sparse: + assert sp.issparse(x) + return x.toarray() + return x + + +@pytest.mark.parametrize( + "sparse_output, p, n", + [ + (False, 1, 0), + (False, 10, 0), + (False, 0, -10), + (False, 1, -1), + (True, 1, 0), + (True, 10, 0), + ], +) +def test_label_binarizer_one_class(sparse_output, p, n): + y = np.array(["a", "a", "a"]) + lb = LabelBinarizer(sparse_output=sparse_output, pos_label=p, neg_label=n) + + # fit_transform + sol = np.array([n, n, n])[:, None] + res = toarray(lb.fit_transform(y), sparse_output) + np.testing.assert_array_equal(res, sol) + + # transform + res = toarray(lb.transform(y), sparse_output) + np.testing.assert_array_equal(res, sol) + + # attributes + assert lb.y_type_ == "binary" + assert not lb.sparse_input_ + np.testing.assert_array_equal(lb.classes_, np.array(["a"])) + + # inverse_transform + np.testing.assert_array_equal(lb.inverse_transform(res), y) + + # transform unseen labels + unseen = np.array(["a", "b"]) + res = toarray(lb.transform(unseen), sparse_output) + sol = np.array([n, n])[:, None] + np.testing.assert_array_equal(res, sol) + + @pytest.mark.parametrize( - "labels", + "sparse_output, p, n", [ - ([1, 4, 5, 2, 0, 1, 6, 2, 3, 4], [4, 2, 6, 3, 2, 0, 1]), - ([9, 8, 2, 1, 3, 4], [8, 2, 1, 2, 2]), + (False, 1, 0), + (False, 10, 0), + (False, 0, -10), + (False, 1, -1), + (True, 1, 0), + (True, 10, 0), ], ) -@pytest.mark.parametrize("dtype", [cp.int32, cp.int64]) -@pytest.mark.parametrize("sparse_output", [True, False]) -def test_basic_functions(labels, dtype, sparse_output): - fit_labels, xform_labels = labels +def test_label_binarizer_two_classes(sparse_output, p, n): + y = np.array(["a", "b", "b", "a"]) + lb = LabelBinarizer(sparse_output=sparse_output, pos_label=p, neg_label=n) + + # fit_transform + sol = np.array([n, p, p, n])[:, None] + res = toarray(lb.fit_transform(y), sparse_output) + np.testing.assert_array_equal(res, sol) + + # transform + res = toarray(lb.transform(y), sparse_output) + np.testing.assert_array_equal(res, sol) + + # attributes + assert lb.y_type_ == "binary" + assert not lb.sparse_input_ + np.testing.assert_array_equal(lb.classes_, np.array(["a", "b"])) + + # inverse_transform + np.testing.assert_array_equal(lb.inverse_transform(res), y) + # can also invert 2 column output + y2 = np.array([[p, n], [n, p], [n, p], [p, n]]) + np.testing.assert_array_equal(lb.inverse_transform(y2), y) + + # transform of unseen classes results in 2 columns + unseen = np.array(["a", "b", "c"]) + res = toarray(lb.transform(unseen), sparse_output) + sol = np.array([[p, n], [n, p], [n, n]]) + np.testing.assert_array_equal(res, sol) + + +@pytest.mark.parametrize( + "sparse_output, p, n", + [ + (False, 1, 0), + (False, 10, 0), + (False, 0, -10), + (False, 1, -1), + (True, 1, 0), + (True, 10, 0), + ], +) +def test_label_binarizer_multiclass(sparse_output, p, n): + y = np.array(["a", "b", "b", "a", "c"]) + lb = LabelBinarizer(sparse_output=sparse_output, pos_label=p, neg_label=n) + + # fit_transform + sol = np.array([[p, n, n], [n, p, n], [n, p, n], [p, n, n], [n, n, p]]) + res = toarray(lb.fit_transform(y), sparse_output) + np.testing.assert_array_equal(res, sol) + + # transform + res = toarray(lb.transform(y), sparse_output) + np.testing.assert_array_equal(res, sol) + + # attributes + assert lb.y_type_ == "multiclass" + assert not lb.sparse_input_ + np.testing.assert_array_equal(lb.classes_, np.array(["a", "b", "c"])) + + # inverse_transform + np.testing.assert_array_equal(lb.inverse_transform(res), y) + + # transform of unseen inputs results in all 0s + unseen = np.array(["d", "a", "e", "c", "f"]) + res = toarray(lb.transform(unseen), sparse_output) + sol = np.array([[n, n, n], [p, n, n], [n, n, n], [n, n, p], [n, n, n]]) + np.testing.assert_array_equal(res, sol) + + +@pytest.mark.parametrize( + "sparse_output, p, n", + [ + (False, 1, 0), + (False, 10, 0), + (False, 0, -10), + (False, 1, -1), + (True, 1, 0), + (True, 10, 0), + ], +) +@pytest.mark.parametrize("sparse_input", [False, True]) +def test_label_binarizer_multilabel_indicator( + sparse_input, sparse_output, p, n +): + y = np.array([[0, 0, 1], [1, 0, 0], [0, 1, 0]]) + if sparse_input: + y = sp.csr_matrix(y) + lb = LabelBinarizer(sparse_output=sparse_output, pos_label=p, neg_label=n) + + # fit_transform + sol = np.array([[n, n, p], [p, n, n], [n, p, n]]) + res = toarray(lb.fit_transform(y), sparse_output) + np.testing.assert_array_equal(res, sol) - skl_bin = skLB(sparse_output=sparse_output) - skl_bin.fit(fit_labels) + # transform + res = toarray(lb.transform(y), sparse_output) + np.testing.assert_array_equal(res, sol) - fit_labels = cp.asarray(fit_labels, dtype=dtype) - xform_labels = cp.asarray(xform_labels, dtype=dtype) + # attributes + assert lb.y_type_ == "multilabel-indicator" + assert lb.sparse_input_ is sparse_input + np.testing.assert_array_equal(lb.classes_, np.array([0, 1, 2])) - binarizer = LabelBinarizer(sparse_output=sparse_output) - binarizer.fit(fit_labels) + # inverse_transform accepts both sparse and dense, + # and returns the type used for fit + sol = np.array([[0, 1, 0], [1, 0, 0]]) + y_dense = np.array([[n, p, n], [p, n, n]]) + y_sparse = sp.csr_matrix(y_dense) + res = toarray(lb.inverse_transform(y_dense), sparse_input) + np.testing.assert_array_equal(res, sol) + res = toarray(lb.inverse_transform(y_sparse), sparse_input) + np.testing.assert_array_equal(res, sol) - assert array_equal(binarizer.classes_.get(), np.unique(fit_labels.get())) + # Can also transform non-multilabel input + y = np.array([3, 2, 1, 0]) + sol = np.array([[n, n, n], [n, n, p], [n, p, n], [p, n, n]]) + res = toarray(lb.transform(y), sparse_output) + np.testing.assert_array_equal(res, sol) - xformed = binarizer.transform(xform_labels) - if sparse_output: - skl_bin_xformed = skl_bin.transform(xform_labels.get()) +@pytest.mark.parametrize("threshold", [1, -1]) +@pytest.mark.parametrize("sparse", [False, True]) +def test_label_binarizer_inverse_transform_threshold(threshold, sparse): + lb = LabelBinarizer().fit(np.array(["a", "b", "c"])) - skl_csr = scipy.sparse.coo_matrix(skl_bin_xformed).tocsr() - cuml_csr = xformed + p = threshold + 1 + n = threshold - 1 + y = np.array([[n, p, n], [p, n, n], [n, n, p]]) + sol = np.array(["b", "a", "c"]) + if sparse: + y = sp.csr_matrix(y) - array_equal(skl_csr.data, cuml_csr.data.get()) + res = lb.inverse_transform(y, threshold=threshold) + np.testing.assert_array_equal(res, sol) - # #todo: Support sparse inputs - # xformed = xformed.todense().astype(dtype) - assert xformed.shape[1] == binarizer.classes_.shape[0] +def test_label_binarize_respects_class_order(): + out = label_binarize([1, 6], classes=[1, 2, 4, 6]) + expected = cp.array([[1, 0, 0, 0], [0, 0, 0, 1]]) + cp.testing.assert_array_equal(out, expected) - original = binarizer.inverse_transform(xformed) + # Modified class order + out = label_binarize([1, 6], classes=[1, 6, 4, 2]) + expected = cp.array([[1, 0, 0, 0], [0, 1, 0, 0]]) + cp.testing.assert_array_equal(out, expected) - assert array_equal(original.get(), xform_labels.get()) + out = label_binarize([0, 1, 2, 3], classes=[3, 2, 0, 1]) + expected = cp.array( + [[0, 0, 1, 0], [0, 0, 0, 1], [0, 1, 0, 0], [1, 0, 0, 0]] + ) + cp.testing.assert_array_equal(out, expected) diff --git a/python/cuml/tests/test_naive_bayes.py b/python/cuml/tests/test_naive_bayes.py index f27561080c..0fd08d1ba8 100644 --- a/python/cuml/tests/test_naive_bayes.py +++ b/python/cuml/tests/test_naive_bayes.py @@ -34,24 +34,17 @@ @pytest.mark.parametrize("x_dtype", [cp.int32, cp.int64]) @pytest.mark.parametrize("y_dtype", [cp.int32, cp.int64]) -def test_sparse_integral_dtype_fails(x_dtype, y_dtype, sparse_text_dataset): +def test_sparse_integral_dtypes_work(x_dtype, y_dtype, sparse_text_dataset): + """These require a bit of special casing since cupyx doesn't support + integral spasre matrices""" X, y = sparse_text_dataset X = X.astype(x_dtype) y = y.astype(y_dtype) - model = MultinomialNB() - - with pytest.raises(ValueError): - model.fit(X, y) - - X = X.astype(cp.float32) - model.fit(X, y) - - X = X.astype(x_dtype) - - with pytest.raises(ValueError): - model.predict(X) + model = MultinomialNB().fit(X, y) + out = model.predict(X) + assert out.dtype == y.dtype @pytest.mark.parametrize("x_dtype", [cp.float32, cp.float64, cp.int32]) diff --git a/python/cuml/tests/test_sklearn_import_export.py b/python/cuml/tests/test_sklearn_import_export.py index c9a0d09319..b49d4a8cf6 100644 --- a/python/cuml/tests/test_sklearn_import_export.py +++ b/python/cuml/tests/test_sklearn_import_export.py @@ -1044,3 +1044,27 @@ def test_label_encoder(): np.testing.assert_array_equal(cu_out, sol) np.testing.assert_array_equal(sk_out, sol) + + +def test_label_binarizer(): + y = np.array(["a", "b", "c", "a"]) + cu_model = cuml.preprocessing.LabelBinarizer().fit(y) + sk_model = sklearn.preprocessing.LabelBinarizer().fit(y) + + cu_model2 = cuml.preprocessing.LabelBinarizer.from_sklearn(sk_model) + sk_model2 = cu_model.as_sklearn() + + roundtrip = cuml.preprocessing.LabelBinarizer.from_sklearn(sk_model2) + assert_roundtrip_consistency(cu_model, roundtrip) + + np.testing.assert_array_equal(cu_model.classes_, sk_model2.classes_) + np.testing.assert_array_equal(cu_model.classes_, roundtrip.classes_) + assert roundtrip.y_type_ == cu_model.y_type_ + assert roundtrip.sparse_input_ == cu_model.sparse_input_ + + sol = np.array([[1, 0, 0], [0, 1, 0], [0, 0, 1], [1, 0, 0]]) + cu_out = cu_model2.transform(y) + sk_out = sk_model2.transform(y) + + np.testing.assert_array_equal(cu_out, sol) + np.testing.assert_array_equal(sk_out, sol) diff --git a/python/cuml/tests/test_validation.py b/python/cuml/tests/test_validation.py index 95a97ca113..bcad6f55e4 100644 --- a/python/cuml/tests/test_validation.py +++ b/python/cuml/tests/test_validation.py @@ -950,15 +950,14 @@ def test_check_array_sparse_not_supported(array): def test_check_array_sparse_input(array, mem_type): ns = cp_sp if is_cuda_output(mem_type, array) else sp # dtype=None case - if ( - mem_type == "device" - and array.dtype.kind != "f" - and array.format != "dia" - ): - # cupy only supports floating dtypes for these inputs. We let cupy - # itself raise an exception, we don't really care what it is. - with pytest.raises(ValueError, match="float32"): - check_array(array, accept_sparse=True, mem_type=mem_type) + if mem_type == "device" and array.dtype.kind != "f": + # coercing to device causes a dtype conversion even if dtype=None if + # the input is an unsupported type. In that case we chose the float + # type with the nearest itemsize. + out = check_array(array, accept_sparse=True, mem_type=mem_type) + coerced_dtype = "float32" if array.dtype.itemsize <= 4 else "float64" + assert ns.issparse(out) + assert out.dtype == coerced_dtype else: out = check_array(array, accept_sparse=True, mem_type=mem_type) assert ns.issparse(out) @@ -972,6 +971,47 @@ def test_check_array_sparse_input(array, mem_type): assert out.dtype == "float32" +def test_check_array_sparse_input_unsupported_dtype(): + """Cupy doesn't support integral types while scipy does, this checks a few + edge cases for how we handle that""" + host_i2 = sp.rand(100, 100, density=0.5).astype("i2") + host_i4 = sp.rand(100, 100, density=0.5).astype("i4") + host_i8 = sp.rand(100, 100, density=0.5).astype("i8") + device_f4 = cp_sp.rand(100, 100, density=0.5).astype("f4") + + # Unsupported dtypes are coerced to floats by default + for array in [host_i2, host_i4, host_i8]: + out = check_array(array, accept_sparse=True) + dtype = "f4" if array.dtype.itemsize <= 4 else "f8" + assert out.dtype == dtype + + # Unsupported dtypes skipped when selecting conversion + out = check_array(host_i4, dtype=["i4", "f8", "f4"], accept_sparse=True) + assert out.dtype == "f8" + + # When coercing host->host, ints are supported + out = check_array(host_i4, accept_sparse=True, mem_type=None) + assert out.dtype == "i4" + out = check_array(host_i2, dtype=["i4"], accept_sparse=True, mem_type=None) + assert out.dtype == "i4" + + # When coercing device->host, ints are supported + out = check_array( + device_f4, dtype=["i4", "i8"], mem_type="host", accept_sparse=True + ) + assert out.dtype == "i4" + + # When coercing to device, error if only unsupported dtypes specified + with pytest.raises(ValueError, match="is supported by cupyx.scipy.sparse"): + check_array(host_i4, dtype=["i4", "i8"], accept_sparse=True) + with pytest.raises(ValueError, match="is supported by cupyx.scipy.sparse"): + check_array(device_f4, dtype=["i4", "i8"], accept_sparse=True) + with pytest.raises(ValueError, match="is supported by cupyx.scipy.sparse"): + check_array( + device_f4, mem_type=None, dtype=["i4", "i8"], accept_sparse=True + ) + + @example( array=cp_sp.csr_matrix(cp.array([[1.0, 0], [0, 0]])), mem_type="host", From 1df5fa929b4e2ce3aeb94e52c6a4b64926593256 Mon Sep 17 00:00:00 2001 From: Steve Collins Date: Tue, 19 May 2026 16:27:41 -0600 Subject: [PATCH 12/17] Deprecate `probability` parameter on `SVC` and `LinearSVC` (#8089) Closes #7982 Mirrors sklearn PR #32050 on cuml.SVC and cuml.LinearSVC. Sentinel + _effective_X property pattern, same as PR #7958. FutureWarning fires from fit when the user passes an explicit value. Also fixes a latent bug in the accel proxy where _gpu_fit was reading self.probability truthily. Default SVC() on small data would have routed through the probability code path since the sentinel string is truthy. Authors: - Steve Collins (https://github.com/switch527) - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Jim Crist-Harif (https://github.com/jcrist) URL: https://github.com/rapidsai/cuml/pull/8089 --- docs/source/cuml-accel/limitations.rst | 6 +- .../cuml/cuml/accel/_overrides/sklearn/svm.py | 18 +-- python/cuml/cuml/svm/linear_svc.py | 32 ++++- python/cuml/cuml/svm/svc.py | 59 ++++++-- python/cuml/cuml/svm/svm_base.pyx | 9 +- .../cuml_accel_tests/integration/test_svc.py | 8 +- .../cuml/cuml_accel_tests/upstream/pytest.ini | 6 +- .../explainer/test_explainer_kernel_shap.py | 4 + .../test_explainer_permutation_shap.py | 4 + python/cuml/tests/test_base.py | 4 + python/cuml/tests/test_linear_svm.py | 4 + python/cuml/tests/test_pickle.py | 14 +- .../cuml/tests/test_sklearn_import_export.py | 54 ++++--- python/cuml/tests/test_svm.py | 136 ++++++++++++++++++ 14 files changed, 293 insertions(+), 65 deletions(-) diff --git a/docs/source/cuml-accel/limitations.rst b/docs/source/cuml-accel/limitations.rst index 5fa70dc860..8a0c0aca3a 100644 --- a/docs/source/cuml-accel/limitations.rst +++ b/docs/source/cuml-accel/limitations.rst @@ -507,7 +507,11 @@ SVC - If ``kernel="precomputed"`` or is a callable. - If ``y`` is multiclass. -- If ``probability=True`` and ``y`` doesn't have at least 5 samples per class. +- If ``probability=True``. The ``probability`` parameter is deprecated in + ``scikit-learn>=1.9``, as well as in ``cuml>=26.06``. We recommend using + wrapping ``SVC`` with ``sklearn.calibration.CalibratedClassifierCV`` like + ``CalibratedClassifierCV(SVC(), ensemble=False)`` instead. This will be + supported across ``scikit-learn`` versions, and won't require CPU fallback. Additional notes: diff --git a/python/cuml/cuml/accel/_overrides/sklearn/svm.py b/python/cuml/cuml/accel/_overrides/sklearn/svm.py index 8798a232af..dff4c89b61 100644 --- a/python/cuml/cuml/accel/_overrides/sklearn/svm.py +++ b/python/cuml/cuml/accel/_overrides/sklearn/svm.py @@ -20,8 +20,8 @@ def _has_probability(model): - # sklearn >= 1.9 defaults `probability` to the sentinel string "deprecated", - # which is truthy. Treat it the same as False (no calibration requested). + # sklearn >= 1.9 defaults `probability` to the "deprecated" sentinel, + # which is truthy. Treat it like False (no calibration requested). if model.probability == "deprecated" or not model.probability: raise AttributeError( "predict_proba is not available when probability=False" @@ -39,26 +39,16 @@ class SVC(ProxyBase): ) def _gpu_fit(self, X, y, sample_weight=None): - classes, counts = np.unique(np.asanyarray(y), return_counts=True) + classes = np.unique(np.asanyarray(y)) if len(classes) > 2: raise UnsupportedOnGPU("Multiclass `y` is not supported") - - # CalibratedClassifierCV doesn't like working with cases where any - # classes have less than 5 examples. - if self.probability and counts.min() < 5: - raise UnsupportedOnGPU( - "`probability=True` requires >= 5 samples per class" - ) - return self._gpu.fit(X, y, sample_weight=sample_weight) def _gpu_decision_function(self, X): # Fixup returned dtype return self._gpu.decision_function(X).astype("float64", copy=False) - # XXX: sklearn wants these methods to only exist if probability=True. - # ProxyBase lacks a builtin mechanism to do that, since this is the only - # use case so far we manually define them for now. + # Manual gate: ProxyBase has no built-in conditional-method support. @available_if(_has_probability) @functools.wraps(_SVC.predict_proba) def predict_proba(self, X): diff --git a/python/cuml/cuml/svm/linear_svc.py b/python/cuml/cuml/svm/linear_svc.py index 689a4637e8..c13f30e9ae 100644 --- a/python/cuml/cuml/svm/linear_svc.py +++ b/python/cuml/cuml/svm/linear_svc.py @@ -1,6 +1,8 @@ # SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # +import warnings + import cupy as cp from sklearn.exceptions import NotFittedError from sklearn.utils.metaestimators import available_if @@ -54,6 +56,11 @@ class LinearSVC(Base, InteropMixin, LinearClassifierMixin, ClassifierMixin): The string 'balanced' is also accepted, in which case ``class_weight[i] = n_samples / (n_classes * n_samples_of_class[i])`` probability: bool, default=False + .. deprecated:: 26.06 + ``probability`` is deprecated and will be removed in version + 26.08. Use ``CalibratedClassifierCV(LinearSVC(), ensemble=False)`` + from ``sklearn.calibration`` for probability estimates instead. + Set to True to enable probability estimate methods (``predict_proba``, ``predict_log_proba``). tol : float, default=1e-4 @@ -155,6 +162,7 @@ def _params_from_cpu(cls, model): f"`multi_class={model.multi_class}` is not supported" ) + # probability omitted: sklearn.LinearSVC has no such param. return { "penalty": model.penalty, "loss": model.loss, @@ -208,7 +216,7 @@ def __init__( fit_intercept=True, penalized_intercept=False, class_weight=None, - probability=False, + probability="deprecated", tol=1e-4, max_iter=1000, linesearch_max_iter=100, @@ -234,12 +242,26 @@ def __init__( self.n_streams = n_streams self.multi_class = multi_class + @property + def _effective_probability(self): + return False if self.probability == "deprecated" else self.probability + @generate_docstring() @reflect(reset="type") def fit( self, X, y, sample_weight=None, *, convert_dtype=True ) -> "LinearSVC": """Fit the model according to the given training data.""" + if self.probability != "deprecated": + warnings.warn( + "The `probability` parameter is deprecated and will be " + "removed in cuML version 26.08. Use " + "`CalibratedClassifierCV(LinearSVC(), ensemble=False)` from " + "`sklearn.calibration` instead.", + FutureWarning, + stacklevel=2, + ) + coef, intercept, n_iter, prob_scale, classes = cuml.svm.linear.fit( self, X, @@ -248,7 +270,7 @@ def fit( convert_dtype=convert_dtype, is_classifier=True, n_streams=self.n_streams, - probability=self.probability, + probability=self._effective_probability, class_weight=self.class_weight, loss=self.loss, penalty=self.penalty, @@ -284,7 +306,7 @@ def fit( @run_in_internal_context def predict(self, X, *, convert_dtype=True): """Predict class labels for samples in X.""" - if self.probability: + if self._effective_probability: scores = self.predict_proba(X, convert_dtype=convert_dtype) else: scores = self.decision_function(X, convert_dtype=convert_dtype) @@ -301,7 +323,7 @@ def predict(self, X, *, convert_dtype=True): inds, self.classes_, output_type=output_type, index=index ) - @available_if(lambda self: self.probability) + @available_if(lambda self: self._effective_probability) @generate_docstring( return_values={ "name": "probs", @@ -329,7 +351,7 @@ def predict_proba(self, X, *, convert_dtype=True) -> CumlArray: n_streams=self.n_streams, ) - @available_if(lambda self: self.probability) + @available_if(lambda self: self._effective_probability) @generate_docstring( return_values={ "name": "probs", diff --git a/python/cuml/cuml/svm/svc.py b/python/cuml/cuml/svm/svc.py index 8160203cf8..bd5295b0ec 100644 --- a/python/cuml/cuml/svm/svc.py +++ b/python/cuml/cuml/svm/svc.py @@ -1,8 +1,12 @@ # SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # +import warnings + import cupy as cp import numpy as np +import sklearn +from packaging.version import Version from sklearn.exceptions import NotFittedError from sklearn.utils.metaestimators import available_if @@ -26,6 +30,8 @@ from cuml.multiclass import OneVsOneClassifier, OneVsRestClassifier from cuml.svm.svm_base import SVMBase +SKLEARN_19 = Version(sklearn.__version__) >= Version("1.9.0.dev0") + class SVC(SVMBase, ClassifierMixin): """ @@ -105,6 +111,11 @@ class SVC(SVMBase, ClassifierMixin): (`cuml.global_settings.output_type`) will be used. See :ref:`output-data-type-configuration` for more info. probability : bool (default = False) + .. deprecated:: 26.06 + ``probability`` is deprecated and will be removed in version + 26.08. Use ``CalibratedClassifierCV(SVC(), ensemble=False)`` + from ``sklearn.calibration`` for probability estimates instead. + Set to ``True`` to enable probability estimates (``predict_proba``/``predict_log_proba``). Note that ``probability=True`` requires your training data have at least 5 @@ -188,18 +199,18 @@ def _get_param_names(cls): @classmethod def _params_from_cpu(cls, model): + if model.probability is True: + # probability=True is deprecated; cuml.accel falls back to + # native sklearn's own CalibratedClassifierCV. + raise UnsupportedOnGPU("`probability=True` is not supported") + params = super()._params_from_cpu(model) params.pop( "epsilon" ) # SVC doesn't expose `epsilon` in the constructor - # sklearn 1.9 changed the default of `probability` from False to the - # sentinel string "deprecated"; coerce to the bool cuml uses. - probability = model.probability - if probability == "deprecated": - probability = False + # probability omitted; True rejected above, default restores it. params.update( { - "probability": probability, "random_state": model.random_state, "class_weight": model.class_weight, "decision_function_shape": model.decision_function_shape, @@ -212,9 +223,14 @@ def _params_to_cpu(self): params.pop( "epsilon" ) # SVC doesn't expose `epsilon` in the constructor + # sklearn <1.9 rejects the ``"deprecated"`` sentinel; resolve it there. + if SKLEARN_19: + probability = self.probability + else: + probability = self._effective_probability params.update( { - "probability": self.probability, + "probability": probability, "random_state": self.random_state, "class_weight": self.class_weight, "decision_function_shape": self.decision_function_shape, @@ -265,7 +281,7 @@ def __init__( nochange_steps=1000, verbose=False, output_type=None, - probability=False, + probability="deprecated", random_state=None, class_weight=None, decision_function_shape="ovo", @@ -288,6 +304,10 @@ def __init__( self.class_weight = class_weight self.decision_function_shape = decision_function_shape + @property + def _effective_probability(self): + return False if self.probability == "deprecated" else self.probability + @property @reflect def support_(self): @@ -326,6 +346,9 @@ def _fit_multiclass(self, X, y, sample_weight): params = self.get_params() decision_function_shape = params.pop("decision_function_shape") + # Pin the sentinel on inner clones so they don't re-fire the + # deprecation warning the outer fit already emitted. + params["probability"] = "deprecated" wrappers = {"ovo": OneVsOneClassifier, "ovr": OneVsRestClassifier} if (multiclass_cls := wrappers.get(decision_function_shape)) is None: raise ValueError( @@ -375,7 +398,7 @@ def _fit_proba(self, X, y, sample_weight): params = { **self.get_params(), - "probability": False, + "probability": "deprecated", "output_type": "numpy", "class_weight": None, } @@ -433,6 +456,16 @@ def fit(self, X, y, sample_weight=None, *, convert_dtype=True) -> "SVC": Fit the model with X and y. """ + if self.probability != "deprecated": + warnings.warn( + "The `probability` parameter is deprecated and will be " + "removed in cuML version 26.08. Use " + "`CalibratedClassifierCV(SVC(), ensemble=False)` from " + "`sklearn.calibration` instead.", + FutureWarning, + stacklevel=2, + ) + if hasattr(self, "_multiclass"): del self._multiclass @@ -476,7 +509,7 @@ def fit(self, X, y, sample_weight=None, *, convert_dtype=True) -> "SVC": balanced_with_sample_weight=False, ) - if self.probability: + if self._effective_probability: return self._fit_proba(X, y, sample_weight) if len(classes) > 2: @@ -507,7 +540,7 @@ def predict(self, X, *, convert_dtype=True): inds = self._multiclass.predict(X) index = inds.index inds = inds.to_output("cupy") - elif self.probability: + elif self._effective_probability: probs = self.predict_proba(X) index = probs.index inds = cp.argmax(probs.to_output("cupy"), axis=1) @@ -522,7 +555,7 @@ def predict(self, X, *, convert_dtype=True): inds, self.classes_, output_type=output_type, index=index ) - @available_if(lambda self: self.probability) + @available_if(lambda self: self._effective_probability) @generate_docstring( skip_parameters_heading=True, return_values={ @@ -584,7 +617,7 @@ def predict_proba(self, X, *, log=False) -> CumlArray: return CumlArray(data=proba, index=index) - @available_if(lambda self: self.probability) + @available_if(lambda self: self._effective_probability) @generate_docstring( return_values={ "name": "preds", diff --git a/python/cuml/cuml/svm/svm_base.pyx b/python/cuml/cuml/svm/svm_base.pyx index c22d9417b9..2f474d5a27 100644 --- a/python/cuml/cuml/svm/svm_base.pyx +++ b/python/cuml/cuml/svm/svm_base.pyx @@ -267,12 +267,9 @@ class SVMBase(Base, "_probA": self._probA, "_probB": self._probB, "_sparse": self._sparse, - # sklearn >= 1.9 added a private fitted attribute that libsvm-facing - # code reads during predict (see sklearn PR #32050). It is set during - # fit(), so we have to set it ourselves. Harmless on older sklearn - # (no code reads it). cuml.SVR has no `probability` attribute; - # default to False to match sklearn. - "_effective_probability": getattr(self, "probability", False), + # sklearn >= 1.9 reads this private attribute during predict + # (sklearn PR #32050); harmless on older sklearn. + "_effective_probability": getattr(self, "_effective_probability", False), **super()._attrs_to_cpu(model), } diff --git a/python/cuml/cuml_accel_tests/integration/test_svc.py b/python/cuml/cuml_accel_tests/integration/test_svc.py index 4f7c480803..642e3463a0 100644 --- a/python/cuml/cuml_accel_tests/integration/test_svc.py +++ b/python/cuml/cuml_accel_tests/integration/test_svc.py @@ -37,14 +37,18 @@ def test_svc(binary): assert svc.score(X, y) > 0.5 +# TODO(26.08): Remove once `probability` is removed from cuml.svm.SVC. @pytest.mark.filterwarnings( - "ignore:The `probability` parameter was deprecated:FutureWarning" + "ignore:Attribute `prob[AB]_` was deprecated:FutureWarning" ) @pytest.mark.filterwarnings( - "ignore:Attribute `prob[AB]_` was deprecated:FutureWarning" + "ignore:The `probability` parameter (is|was) deprecated:FutureWarning" ) def test_svc_probability(binary): X, y = binary + # cuml.accel no longer accelerates `probability=True`; this fit falls + # back to native sklearn `SVC(probability=True)`. predict_proba still + # works through the native path. svc = SVC(probability=True).fit(X, y) # Inference and score works assert svc.score(X, y) > 0.5 diff --git a/python/cuml/cuml_accel_tests/upstream/pytest.ini b/python/cuml/cuml_accel_tests/upstream/pytest.ini index a47cb90c34..c895ea73f2 100644 --- a/python/cuml/cuml_accel_tests/upstream/pytest.ini +++ b/python/cuml/cuml_accel_tests/upstream/pytest.ini @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # pytest config to be used when running the upstream test suites. @@ -14,5 +14,9 @@ filterwarnings = error::cuml.accel.pytest_plugin.UnmatchedXfailTests # Error for FutureWarnings raised by cuml error::FutureWarning:cuml + # Suppress cuml's deprecation warning for sklearn upstream tests that + # legitimately use `probability=True`. Must come after the generic error + # rule (later filterwarnings entries take precedence). TODO(26.08): drop. + ignore:The `probability` parameter is deprecated:FutureWarning # Ignore unknown pytest marks, the xfail-list currently adds a bunch of these ignore::pytest.PytestUnknownMarkWarning diff --git a/python/cuml/tests/explainer/test_explainer_kernel_shap.py b/python/cuml/tests/explainer/test_explainer_kernel_shap.py index 4531c9a279..047df1daa5 100644 --- a/python/cuml/tests/explainer/test_explainer_kernel_shap.py +++ b/python/cuml/tests/explainer/test_explainer_kernel_shap.py @@ -84,6 +84,10 @@ def test_exact_regression_datasets(exact_shap_regression_dataset, model): ) +# TODO(26.08): Remove this filter once `probability` is removed from cuml.svm.SVC. +@pytest.mark.filterwarnings( + "ignore:The `probability` parameter is deprecated:FutureWarning" +) def test_exact_classification_datasets(exact_shap_classification_dataset): X_train, X_test, y_train, y_test = exact_shap_classification_dataset diff --git a/python/cuml/tests/explainer/test_explainer_permutation_shap.py b/python/cuml/tests/explainer/test_explainer_permutation_shap.py index db88336297..b80b44b24a 100644 --- a/python/cuml/tests/explainer/test_explainer_permutation_shap.py +++ b/python/cuml/tests/explainer/test_explainer_permutation_shap.py @@ -54,6 +54,10 @@ def test_regression_datasets(exact_shap_regression_dataset, model): ) <= 1e-5 +# TODO(26.08): Remove this filter once `probability` is removed from cuml.svm.SVC. +@pytest.mark.filterwarnings( + "ignore:The `probability` parameter is deprecated:FutureWarning" +) def test_exact_classification_datasets(exact_shap_classification_dataset): X_train, X_test, y_train, y_test = exact_shap_classification_dataset diff --git a/python/cuml/tests/test_base.py b/python/cuml/tests/test_base.py index c2752f976f..ad440f74dd 100644 --- a/python/cuml/tests/test_base.py +++ b/python/cuml/tests/test_base.py @@ -343,6 +343,10 @@ def test_regressor_predict_dtype(cls): ) # TODO(26.08) Remove this filter @pytest.mark.filterwarnings("ignore:The default value of 'max_depth'") +# TODO(26.08): Remove once `probability` is removed from cuml.svm.SVC. +@pytest.mark.filterwarnings( + "ignore:The `probability` parameter is deprecated:FutureWarning" +) @pytest.mark.parametrize( "target_kind", ["binary", "multiclass", "multitarget"] ) diff --git a/python/cuml/tests/test_linear_svm.py b/python/cuml/tests/test_linear_svm.py index bd2d35e4d5..41182aecb4 100644 --- a/python/cuml/tests/test_linear_svm.py +++ b/python/cuml/tests/test_linear_svm.py @@ -214,6 +214,10 @@ def test_linear_svc_decision_function( @pytest.mark.parametrize("fit_intercept", [True, False]) @pytest.mark.parametrize("n_classes", [2, 3, 5]) +# TODO(26.08): Remove once `probability` is removed from cuml.svm.LinearSVC. +@pytest.mark.filterwarnings( + "ignore:The `probability` parameter is deprecated:FutureWarning" +) def test_linear_svc_predict_proba(fit_intercept, n_classes): n_rows, n_cols = 500, 20 X_train, X_test, y_train, y_test = make_classification_dataset( diff --git a/python/cuml/tests/test_pickle.py b/python/cuml/tests/test_pickle.py index 6c9537f64a..590b6ea4f9 100644 --- a/python/cuml/tests/test_pickle.py +++ b/python/cuml/tests/test_pickle.py @@ -29,10 +29,16 @@ ) from cuml.tsa.arima import ARIMA -# TODO(26.08) Remove this filter -pytestmark = pytest.mark.filterwarnings( - "ignore:The default value of 'max_depth':FutureWarning" -) +pytestmark = [ + # TODO(26.08): Remove this filter + pytest.mark.filterwarnings( + "ignore:The default value of 'max_depth':FutureWarning" + ), + # TODO(26.08): Remove once `probability` is removed from cuml.svm.SVC/LinearSVC. + pytest.mark.filterwarnings( + "ignore:The `probability` parameter is deprecated:FutureWarning" + ), +] regression_config = ClassEnumerator(module=cuml.linear_model) regression_models = regression_config.get_models() diff --git a/python/cuml/tests/test_sklearn_import_export.py b/python/cuml/tests/test_sklearn_import_export.py index b49d4a8cf6..92f2aa780d 100644 --- a/python/cuml/tests/test_sklearn_import_export.py +++ b/python/cuml/tests/test_sklearn_import_export.py @@ -376,8 +376,9 @@ def test_svr(random_state, sparse, kernel): ) +# TODO(26.08): Remove this filter once `probability` is removed from cuml.svm.SVC. @pytest.mark.filterwarnings( - "ignore:The `probability` parameter was deprecated:FutureWarning" + "ignore:The `probability` parameter (is|was) deprecated:FutureWarning" ) @pytest.mark.filterwarnings( "ignore:Attribute `prob[AB]_` was deprecated:FutureWarning" @@ -396,47 +397,62 @@ def test_svc(random_state, sparse, probability, kernel): ) K = X @ X.T # Linear kernel matrix original = cuml.SVC(kernel="precomputed") - assert_estimator_roundtrip(original, sklearn.svm.SVC, K, y) + assert_estimator_roundtrip( + original, + sklearn.svm.SVC, + K, + y, + exclude_params=["probability"], + ) else: if sparse: X = scipy.sparse.coo_matrix(X) original = cuml.SVC() - assert_estimator_roundtrip(original, sklearn.svm.SVC, X, y) + assert_estimator_roundtrip( + original, + sklearn.svm.SVC, + X, + y, + exclude_params=["probability"], + ) # Check inference works after conversion. sklearn 1.9 deprecated the # `probability` parameter; avoid passing it on the sklearn side when # False (the default) to avoid the FutureWarning. cu_model = cuml.SVC(probability=probability).fit(X, y) - if probability: - sk_model = sklearn.svm.SVC(probability=True).fit(X, y) - else: - sk_model = sklearn.svm.SVC().fit(X, y) - - cu_model2 = cuml.SVC.from_sklearn(sk_model) sk_model2 = cu_model.as_sklearn() - - cu_score = cu_model2.score(X, y) - assert cu_score > 0.7 - sk_score = sk_model2.score(X, y) assert sk_score > 0.7 if probability: - # Check that predict_proba works - cu_pred_prob = cu_model2.predict_proba(X).argmax(axis=1) - assert accuracy_score(cu_pred_prob, y) > 0.7 + # `cuml.SVC.from_sklearn` rejects probability=True so cuml.accel + # falls back to native sklearn for calibrated SVC. + sk_model = sklearn.svm.SVC(probability=True).fit(X, y) + with pytest.raises( + UnsupportedOnGPU, + match=r"`probability=True` is not supported", + ): + cuml.SVC.from_sklearn(sk_model) + + # The cuml-direct path (used here via cu_model.as_sklearn()) still + # wires up probA_ / probB_ for predict_proba on the sklearn side. sk_pred_prob = sk_model2.predict_proba(X).argmax(axis=1) assert accuracy_score(sk_pred_prob, y) > 0.7 - # Check that probA_, probB_ are wired up properly for attr in ["probA_", "probB_"]: val = getattr(sk_model2, attr) assert isinstance(val, np.ndarray) assert val.dtype == "float64" assert val.shape == (1,) + else: + sk_model = sklearn.svm.SVC().fit(X, y) + cu_model2 = cuml.SVC.from_sklearn(sk_model) + + cu_score = cu_model2.score(X, y) + assert cu_score > 0.7 - # Check n_support_ is correctly set - assert cu_model2.n_support_ == cu_model2.support_vectors_.shape[0] + # Check n_support_ is correctly set + assert cu_model2.n_support_ == cu_model2.support_vectors_.shape[0] @pytest.mark.parametrize("kind", ["SVC", "SVR"]) diff --git a/python/cuml/tests/test_svm.py b/python/cuml/tests/test_svm.py index 207d218723..107c7e622d 100644 --- a/python/cuml/tests/test_svm.py +++ b/python/cuml/tests/test_svm.py @@ -2,6 +2,7 @@ # SPDX-License-Identifier: Apache-2.0 # import platform +import warnings import cudf import cupy as cp @@ -37,6 +38,13 @@ unit_param, ) +# Many tests below pass `probability=` to cuml SVC/LinearSVC on purpose; +# silence the FutureWarning module-wide. +# TODO(26.08): Remove once `probability` is removed from cuml.svm.SVC/LinearSVC. +pytestmark = pytest.mark.filterwarnings( + "ignore:The `probability` parameter is deprecated:FutureWarning" +) + IS_ARM = platform.processor() == "aarch64" @@ -645,6 +653,134 @@ def test_svc_probability_n_iter(): assert model.n_iter_.shape == (1,) +@pytest.mark.parametrize("explicit_value", [True, False]) +def test_svc_probability_emits_future_warning(explicit_value): + # cuml-side half of #7982. + X, y = make_classification( + n_samples=40, + n_features=4, + n_informative=3, + n_redundant=0, + n_classes=2, + random_state=0, + ) + with pytest.warns( + FutureWarning, + match=( + r"The `probability` parameter is deprecated and will be " + r"removed in cuML" + ), + ): + cu_svm.SVC(probability=explicit_value).fit(X, y) + + +def test_svc_default_does_not_emit_probability_warning(): + # If the user does not pass `probability=`, no FutureWarning should fire. + # The default sentinel string `"deprecated"` is treated as `probability=False`. + X, y = make_classification( + n_samples=40, + n_features=4, + n_informative=3, + n_redundant=0, + n_classes=2, + random_state=0, + ) + with warnings.catch_warnings(): + warnings.filterwarnings( + "error", + message=( + r"The `probability` parameter is deprecated and will " + r"be removed in cuML" + ), + category=FutureWarning, + ) + cu_svm.SVC().fit(X, y) + + +@pytest.mark.parametrize("explicit_value", [True, False]) +def test_linear_svc_probability_emits_future_warning(explicit_value): + X, y = make_classification( + n_samples=40, + n_features=4, + n_informative=3, + n_redundant=0, + n_classes=2, + random_state=0, + ) + with pytest.warns( + FutureWarning, + match=( + r"The `probability` parameter is deprecated and will be " + r"removed in cuML" + ), + ): + cu_svm.LinearSVC(probability=explicit_value).fit(X, y) + + +def test_linear_svc_default_does_not_emit_probability_warning(): + X, y = make_classification( + n_samples=40, + n_features=4, + n_informative=3, + n_redundant=0, + n_classes=2, + random_state=0, + ) + with warnings.catch_warnings(): + warnings.filterwarnings( + "error", + message=( + r"The `probability` parameter is deprecated and will " + r"be removed in cuML" + ), + category=FutureWarning, + ) + cu_svm.LinearSVC().fit(X, y) + + +def test_svc_probability_warning_fires_once(): + # Inner clones (CalibratedClassifierCV folds, multiclass binaries) pin + # the sentinel so the FutureWarning fires only on the outer fit, not + # per inner fit. + X_bin, y_bin = make_classification( + n_samples=40, + n_features=4, + n_informative=3, + n_redundant=0, + n_classes=2, + random_state=0, + ) + with warnings.catch_warnings(record=True) as caught: + warnings.simplefilter("always") + cu_svm.SVC(probability=True).fit(X_bin, y_bin) + matched = [ + w + for w in caught + if issubclass(w.category, FutureWarning) + and "is deprecated and will be removed in cuML" in str(w.message) + ] + assert len(matched) == 1 + + X_mc, y_mc = make_classification( + n_samples=60, + n_features=4, + n_informative=3, + n_redundant=0, + n_classes=3, + random_state=0, + ) + with warnings.catch_warnings(record=True) as caught: + warnings.simplefilter("always") + cu_svm.SVC(probability=False).fit(X_mc, y_mc) + matched = [ + w + for w in caught + if issubclass(w.category, FutureWarning) + and "is deprecated and will be removed in cuML" in str(w.message) + ] + assert len(matched) == 1 + + # Tests for kernel='precomputed' @pytest.mark.parametrize("kernel_func", ["linear", "rbf"]) def test_svc_precomputed_kernel(kernel_func): From 2e562d1d97ea3efa9ce41bed57da1dd3ca41c9d5 Mon Sep 17 00:00:00 2001 From: Jim Crist-Harif Date: Wed, 20 May 2026 22:50:15 -0500 Subject: [PATCH 13/17] Further relax `test_nearest_neighbors_pickle` tolerance (#8136) This failure started after the recent CCCL upgrade and appears only on rtxpro6000 test runs. We relaxed the tolerance once already, but have still seen a rare periodic failure. From looking at recent failures, an atol of 5e-3 would be sufficient, but bumping to 1e-2 to be sure. We're only checking plumbing here, so some slop in tolerance is fine. Authors: - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Simon Adorf (https://github.com/csadorf) - Victor Lafargue (https://github.com/viclafargue) URL: https://github.com/rapidsai/cuml/pull/8136 --- python/cuml/tests/test_pickle.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/cuml/tests/test_pickle.py b/python/cuml/tests/test_pickle.py index 590b6ea4f9..beb0973a70 100644 --- a/python/cuml/tests/test_pickle.py +++ b/python/cuml/tests/test_pickle.py @@ -479,7 +479,7 @@ def test_nearest_neighbors_pickle(algorithm): # just to ensure things are wired together properly. accuracy = (i1 == i2).sum() / i1.size assert accuracy >= 0.9 - np.testing.assert_allclose(d1, d2, atol=1e-3) + np.testing.assert_allclose(d1, d2, atol=1e-2) else: np.testing.assert_allclose(i1, i2) np.testing.assert_allclose(d1, d2) From 651e240e9b52d28b8b3315a82a99e202016a06f1 Mon Sep 17 00:00:00 2001 From: John Zedlewski <904524+JohnZed@users.noreply.github.com> Date: Wed, 20 May 2026 22:14:51 -0700 Subject: [PATCH 14/17] Add cuml.accel support for IncrementalPCA (#7785) Adds `sklearn.decomposition.IncrementalPCA` to `cuml.accel`, including `fit_transform` and `partial_fit` dispatch through the estimator proxy. This also tightens sklearn parity for cuML `IncrementalPCA` fitted attributes, CPU/GPU interop, `var_` handling, first-batch validation behavior, and `set_output` support. Closes #7779 Authors: - John Zedlewski (https://github.com/JohnZed) - Simon Adorf (https://github.com/csadorf) Approvers: - Simon Adorf (https://github.com/csadorf) URL: https://github.com/rapidsai/cuml/pull/7785 --- docs/source/cuml-accel/faq.rst | 1 + docs/source/cuml-accel/limitations.rst | 10 +++ .../accel/_overrides/sklearn/decomposition.py | 20 ++++- .../cuml/decomposition/incremental_pca.py | 80 +++++++++++++++++-- .../cuml_accel_tests/test_estimator_proxy.py | 53 +++++++++++- python/cuml/tests/test_incremental_pca.py | 49 ++++++++++++ 6 files changed, 204 insertions(+), 9 deletions(-) diff --git a/docs/source/cuml-accel/faq.rst b/docs/source/cuml-accel/faq.rst index e83b724749..1b3f925bd7 100644 --- a/docs/source/cuml-accel/faq.rst +++ b/docs/source/cuml-accel/faq.rst @@ -55,6 +55,7 @@ the following estimators are mostly or entirely accelerated when run with * ``sklearn.covariance.EmpiricalCovariance`` * ``sklearn.covariance.LedoitWolf`` * ``sklearn.decomposition.PCA`` + * ``sklearn.decomposition.IncrementalPCA`` * ``sklearn.decomposition.TruncatedSVD`` * ``sklearn.ensemble.RandomForestClassifier`` * ``sklearn.ensemble.RandomForestRegressor`` diff --git a/docs/source/cuml-accel/limitations.rst b/docs/source/cuml-accel/limitations.rst index 8a0c0aca3a..9a35f94cac 100644 --- a/docs/source/cuml-accel/limitations.rst +++ b/docs/source/cuml-accel/limitations.rst @@ -156,6 +156,16 @@ Additional notes: - Parameters for the ``"randomized"`` solver like ``random_state``, ``n_oversamples``, ``power_iteration_normalizer`` are ignored. +IncrementalPCA +^^^^^^^^^^^^^^ + +``IncrementalPCA`` has no known estimator-specific ``cuml.accel`` limitations. + +Additional notes: + +- ``partial_fit`` does not support sparse input. This matches scikit-learn; + use ``fit`` for sparse input or provide dense batches to ``partial_fit``. + TruncatedSVD ^^^^^^^^^^^^ diff --git a/python/cuml/cuml/accel/_overrides/sklearn/decomposition.py b/python/cuml/cuml/accel/_overrides/sklearn/decomposition.py index adb5ef8e55..026851cb56 100644 --- a/python/cuml/cuml/accel/_overrides/sklearn/decomposition.py +++ b/python/cuml/cuml/accel/_overrides/sklearn/decomposition.py @@ -1,5 +1,5 @@ # -# SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # @@ -9,7 +9,7 @@ import cuml.decomposition from cuml.accel.estimator_proxy import ProxyBase -__all__ = ("PCA", "TruncatedSVD") +__all__ = ("IncrementalPCA", "PCA", "TruncatedSVD") # In sklearn 1.5 the sign flipping behavior changed. For sklearn < 1.5 we # enable the old behavior. @@ -42,3 +42,19 @@ def _gpu_fit(self, X, y=None): def _gpu_fit_transform(self, X, y=None): self._gpu._u_based_sign_flip = True return self._gpu.fit_transform(X, y) + + +class IncrementalPCA(ProxyBase): + _gpu_class = cuml.decomposition.IncrementalPCA + + def _gpu_fit_transform(self, X, y=None, **fit_params): + if fit_params: + param = next(iter(fit_params)) + raise TypeError( + "IncrementalPCA.fit() got an unexpected keyword argument " + f"{param!r}" + ) + return self._gpu.fit_transform(X, y) + + def _gpu_partial_fit(self, X, y=None, check_input=True): + return self._gpu.partial_fit(X, y, check_input=check_input) diff --git a/python/cuml/cuml/decomposition/incremental_pca.py b/python/cuml/cuml/decomposition/incremental_pca.py index b89a1e73ca..1997cf2466 100644 --- a/python/cuml/cuml/decomposition/incremental_pca.py +++ b/python/cuml/cuml/decomposition/incremental_pca.py @@ -8,10 +8,12 @@ import cupy as cp import cuml.internals +from cuml.common.array_descriptor import CumlArrayDescriptor from cuml.common.sparse_utils import is_sparse from cuml.decomposition.pca import PCA from cuml.internals.array import CumlArray from cuml.internals.base import Base +from cuml.internals.interop import InteropMixin, to_cpu, to_gpu from cuml.internals.validation import ( check_array, check_features, @@ -177,6 +179,9 @@ class IncrementalPCA(PCA): 0.0037122774558343763 """ + _cpu_class_path = "sklearn.decomposition.IncrementalPCA" + var_ = CumlArrayDescriptor(order="F") + def __init__( self, *, @@ -279,7 +284,10 @@ def partial_fit(self, X, y=None, *, check_input=True) -> "IncrementalPCA": if first_call := getattr(self, "n_samples_seen_", 0) == 0: self._set_output_type(X) - check_features(self, X, reset=first_call) + # `fit()` pre-validates X once and calls us per batch with + # check_input=False; skip re-running check_features in that case. + if check_input or not hasattr(self, "n_features_in_"): + check_features(self, X, reset=first_call) if check_input: X = check_array(X, dtype=("float32", "float64")) @@ -309,11 +317,11 @@ def partial_fit(self, X, y=None, *, check_input=True) -> "IncrementalPCA": "more rows than columns for IncrementalPCA " "processing" % (self.n_components, n_features) ) - elif not self.n_components <= n_samples: + elif self.n_components > n_samples and first_call: raise ValueError( - "n_components=%r must be less or equal to " - "the batch number of samples " - "%d." % (self.n_components, n_samples) + f"n_components={self.n_components} must be less or equal to " + f"the batch number of samples {n_samples} for the first " + "partial_fit call." ) else: self.n_components_ = self.n_components @@ -371,7 +379,7 @@ def partial_fit(self, X, y=None, *, check_input=True) -> "IncrementalPCA": ) self.singular_values_ = CumlArray(data=S[: self.n_components_]) self.mean_ = CumlArray(data=col_mean) - self.var_ = col_var + self.var_ = CumlArray(data=col_var) self.explained_variance_ = CumlArray( data=explained_variance[: self.n_components_] ) @@ -454,6 +462,66 @@ def _get_param_names(cls): "batch_size", ] + @classmethod + def _params_from_cpu(cls, model): + return { + "n_components": model.n_components, + "whiten": model.whiten, + "copy": model.copy, + "batch_size": model.batch_size, + } + + def _params_to_cpu(self): + return { + "n_components": self.n_components, + "whiten": self.whiten, + "copy": self.copy, + "batch_size": self.batch_size, + } + + # Bypass PCA._attrs_*: sklearn IncrementalPCA lacks PCA's `n_samples_`. + # Call InteropMixin directly for universal `n_features_in_` / + # `feature_names_in_` handling. + def _attrs_from_cpu(self, model): + out = { + "components_": to_gpu(model.components_, order="F"), + "explained_variance_": to_gpu( + model.explained_variance_, order="F" + ), + "explained_variance_ratio_": to_gpu( + model.explained_variance_ratio_, order="F" + ), + "singular_values_": to_gpu(model.singular_values_, order="F"), + "mean_": to_gpu(model.mean_, order="F"), + "var_": to_gpu(model.var_, order="F"), + "n_components_": model.n_components_, + "n_samples_seen_": model.n_samples_seen_, + "noise_variance_": model.noise_variance_, + **InteropMixin._attrs_from_cpu(self, model), + } + if (batch_size_ := getattr(model, "batch_size_", None)) is not None: + out["batch_size_"] = batch_size_ + return out + + def _attrs_to_cpu(self, model): + out = { + "components_": to_cpu(self.components_), + "explained_variance_": to_cpu(self.explained_variance_), + "explained_variance_ratio_": to_cpu( + self.explained_variance_ratio_ + ), + "singular_values_": to_cpu(self.singular_values_), + "mean_": to_cpu(self.mean_), + "var_": to_cpu(self.var_), + "n_components_": self.n_components_, + "n_samples_seen_": self.n_samples_seen_, + "noise_variance_": self.noise_variance_, + **InteropMixin._attrs_to_cpu(self, model), + } + if (batch_size_ := getattr(self, "batch_size_", None)) is not None: + out["batch_size_"] = batch_size_ + return out + def _gen_batches(n, batch_size, min_batch_size=0): """ diff --git a/python/cuml/cuml_accel_tests/test_estimator_proxy.py b/python/cuml/cuml_accel_tests/test_estimator_proxy.py index f35d3d9ff2..cca1f9a6e1 100644 --- a/python/cuml/cuml_accel_tests/test_estimator_proxy.py +++ b/python/cuml/cuml_accel_tests/test_estimator_proxy.py @@ -17,7 +17,7 @@ from packaging.version import Version from sklearn.base import check_is_fitted, is_classifier, is_regressor from sklearn.datasets import make_blobs, make_classification, make_regression -from sklearn.decomposition import PCA +from sklearn.decomposition import PCA, IncrementalPCA from sklearn.ensemble import RandomForestRegressor from sklearn.exceptions import NotFittedError from sklearn.linear_model import ( @@ -142,6 +142,15 @@ def test_init_positional_and_keyword(): # Can't pass keyword-only parameters in as positional PCA(10, False) + model = IncrementalPCA() + assert model.n_components is None + + model = IncrementalPCA(n_components=10) + assert model.n_components == 10 + + model = IncrementalPCA(10) + assert model.n_components == 10 + def test_repr(): model = LogisticRegression(C=1.5) @@ -697,6 +706,48 @@ def test_set_output(methods): assert not hasattr(model._cpu, "n_features_in_") +def test_incremental_pca_partial_fit(): + X, _ = make_classification(n_samples=40, n_features=8, random_state=42) + model = IncrementalPCA(n_components=3, batch_size=10) + + model.partial_fit(X[:20]) + gpu = model._gpu + + assert gpu is not None + assert int(gpu.n_samples_seen_) == 20 + + model.partial_fit(X[20:]) + + assert model._gpu is gpu + assert int(model._gpu.n_samples_seen_) == 40 + assert not hasattr(model._cpu, "components_") + + out = model.transform(X) + assert isinstance(out, np.ndarray) + assert out.shape[1] == 3 + + +def test_incremental_pca_set_output(): + X, _ = make_classification(n_samples=40, n_features=8, random_state=42) + model = IncrementalPCA(n_components=3, batch_size=10) + + assert model.set_output(transform="pandas") is model + model.partial_fit(X) + + out = model.transform(X) + assert isinstance(out, pd.DataFrame) + assert out.columns[0].startswith("incrementalpca") + + +@pytest.mark.parametrize("method", ["fit_transform", "partial_fit"]) +def test_incremental_pca_rejects_unsupported_fit_params(method): + X, _ = make_classification(n_samples=40, n_features=8, random_state=42) + model = IncrementalPCA(n_components=3, batch_size=10) + + with pytest.raises(TypeError): + getattr(model, method)(X, sample_weight=np.ones(X.shape[0])) + + def test_get_feature_names_out(): model = PCA(n_components=5) diff --git a/python/cuml/tests/test_incremental_pca.py b/python/cuml/tests/test_incremental_pca.py index 021e0dd057..99abd6a9dc 100644 --- a/python/cuml/tests/test_incremental_pca.py +++ b/python/cuml/tests/test_incremental_pca.py @@ -112,6 +112,55 @@ def test_partial_fit( assert array_equal(cu_inv, sk_inv, 1e-3, with_sign=True) +@pytest.mark.parametrize("whiten", [False, True]) +@pytest.mark.parametrize("method", ["fit", "partial_fit"]) +def test_sklearn_parity_attrs(method, whiten): + """Check fitted-attribute parity with sklearn for fit() and partial_fit().""" + nrows, ncols, n_components = 1000, 8, 4 + X, _ = make_blobs( + n_samples=nrows, n_features=ncols, random_state=0, dtype="float64" + ) + X_np = cp.asnumpy(X) + + cu_ipca = cuIPCA(n_components=n_components, whiten=whiten, batch_size=100) + sk_ipca = skIPCA(n_components=n_components, whiten=whiten, batch_size=100) + + if method == "fit": + cu_ipca.fit(X) + sk_ipca.fit(X_np) + else: + batch = 100 + for i in range(0, nrows, batch): + cu_ipca.partial_fit(X[i : i + batch]) + sk_ipca.partial_fit(X_np[i : i + batch]) + + assert array_equal( + cu_ipca.components_, sk_ipca.components_, 1e-3, with_sign=True + ) + assert array_equal( + cu_ipca.explained_variance_, + sk_ipca.explained_variance_, + 1e-3, + with_sign=True, + ) + assert array_equal( + cu_ipca.explained_variance_ratio_, + sk_ipca.explained_variance_ratio_, + 1e-3, + with_sign=True, + ) + assert array_equal( + cu_ipca.singular_values_, + sk_ipca.singular_values_, + 1e-3, + with_sign=True, + ) + assert array_equal(cu_ipca.mean_, sk_ipca.mean_, 1e-3, with_sign=True) + assert array_equal(cu_ipca.var_, sk_ipca.var_, 1e-3, with_sign=True) + assert int(cu_ipca.n_samples_seen_) == int(sk_ipca.n_samples_seen_) + assert cu_ipca.n_components_ == sk_ipca.n_components_ + + def test_exceptions(): X = cupyx.scipy.sparse.eye(10) ipca = cuIPCA() From c2eaf61e6751c729d74ca8a1ffe4c8e56cf54035 Mon Sep 17 00:00:00 2001 From: Jim Crist-Harif Date: Thu, 21 May 2026 14:07:39 -0500 Subject: [PATCH 15/17] Add "Advanced Topics" doc (#8134) Adds a new doc with some advanced topics. Includes: - CUDA streams and synchronization - CUDA device selection - A very small section on configuring RMM. I wanted to mention something, but not provide a full guide. Once the upstream RMM one lands we can link there. xref https://github.com/rapidsai/cuml/issues/8128 Authors: - Jim Crist-Harif (https://github.com/jcrist) Approvers: - Gil Forsyth (https://github.com/gforsyth) - Simon Adorf (https://github.com/csadorf) URL: https://github.com/rapidsai/cuml/pull/8134 --- .gitignore | 3 +- docs/source/advanced.rst | 94 ++++++++++++++++++++++++++++++++++++++ docs/source/conf.py | 6 +++ docs/source/user_guide.rst | 1 + 4 files changed, 103 insertions(+), 1 deletion(-) create mode 100644 docs/source/advanced.rst diff --git a/.gitignore b/.gitignore index ce006d831c..bb3ead8e63 100644 --- a/.gitignore +++ b/.gitignore @@ -36,11 +36,12 @@ _skbuild/ junit-*.xml -## files pickled in notebook when ran during python docstring generation +## files generated by notebooks during docs builds docs/source/*.model docs/source/*.pkl docs/source/*.tl docs/source/*.joblib +docs/source/**/*.onnx ## autosummary-generated API reference stubs docs/source/api/generated/ diff --git a/docs/source/advanced.rst b/docs/source/advanced.rst new file mode 100644 index 0000000000..d017924b9f --- /dev/null +++ b/docs/source/advanced.rst @@ -0,0 +1,94 @@ +Advanced Topics +=============== + +Here we cover a few assorted topics that may be of interest to more advanced +use cases. + +CUDA Streams and Synchronization +-------------------------------- + +Functions and methods in cuML are written using a variety of technologies. As +such, while *most* methods run on the CUDA `per-thread default stream`_ (PTDS), +some methods might run on the `legacy default stream`_ (also known as the NULL +stream) instead. + +cuML does not currently expose stream selection as part of its public API and +makes no guarantees on whether a particular method runs on the PTDS or legacy +default stream. Likewise there is no guarantee that the output of a cuML method +or function has been synchronized before returning. + +For users, if you follow the following guideline you shouldn't have any +concurrency issues: + +- Device memory input arrays should be either fully computed, or currently + computing on the PTDS or legacy default stream. + +- Device memory output arrays should be operated on using either the PTDS or + legacy default stream, OR have the PTDS of the thread that ran the method + synchronized before further access. + +- Inputs and outputs using host memory have no restrictions and shouldn't be + prone to concurrency issues. + + +Selecting the CUDA Device +------------------------- + +All single-GPU cuML methods run on device 0 by default. Setting a device via +the :class:`cupy.cuda.Device` or :class:`cuda.core.Device` APIs is currently +not supported. To specify a device to run on, we recommend using the +``CUDA_VISIBLE_DEVICES`` (`doc +`_) +environment variable. For example: + +.. code-block:: shell + + CUDA_VISIBLE_DEVICES=2 python myscript.py + +cuML does contain a few single-node multi-GPU implementations. When available, +these take a ``device_ids`` parameter to specify which devices to run on. See +the `cuml.manifold.UMAP` docs for an example. + + +Configuring the Memory Allocator +-------------------------------- + +Memory allocations in cuML are made using the `Rapids Memory Manager`_ (RMM). +We don't do any configuration of RMM on import; allocations are made using the +default memory resource (:class:`rmm.mr.CudaMemoryResource`). + +Some applications may run better using an alternative memory resource. A few +common options: + +- A good default to try is the :class:`rmm.mr.CudaAsyncMemoryResource`. This is + a stream-ordered pooling resource, and may be faster for your application. + + .. code-block:: python + + import rmm + + rmm.mr.set_current_device_resource(rmm.mr.CudaAsyncMemoryResource()) + + +- Users working with large data may want to enable cuML to use `CUDA Unified + Memory`_ to enable GPU memory oversubscription. To do this, we recommend + using :class:`rmm.mr.ManagedMemoryResource` wrapped in a + :class:`rmm.mr.PrefetchResourceAdaptor` to minimize paging overhead. + + .. code-block:: python + + import rmm + + rmm.mr.set_current_device_resource( + rmm.mr.PrefetchResourceAdaptor(rmm.mr.ManagedMemoryResource()) + ) + + +For more details, see the `RMM documentation`_. + + +.. _per-thread default stream: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/asynchronous-execution.html#per-thread-default-stream +.. _legacy default stream: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/asynchronous-execution.html#legacy-default-stream +.. _Rapids Memory Manager: +.. _RMM documentation: https://docs.rapids.ai/api/rmm/stable/ +.. _CUDA Unified Memory: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/unified-memory.html diff --git a/docs/source/conf.py b/docs/source/conf.py index 97e55066d3..17e1c8d6eb 100644 --- a/docs/source/conf.py +++ b/docs/source/conf.py @@ -203,6 +203,12 @@ # TODO: re-enable once scipy docs are more reliable # "scipy": ("https://docs.scipy.org/doc/scipy", None), "sklearn": ("https://scikit-learn.org/stable/", None), + "cupy": ("https://docs.cupy.dev/en/stable/", None), + "cuda.core": ( + "https://nvidia.github.io/cuda-python/cuda-core/latest/", + None, + ), + "rmm": ("https://docs.rapids.ai/api/rmm/stable/", None), } # Config numpydoc diff --git a/docs/source/user_guide.rst b/docs/source/user_guide.rst index 87a14af895..249f2fb938 100644 --- a/docs/source/user_guide.rst +++ b/docs/source/user_guide.rst @@ -21,5 +21,6 @@ GitHub repository instead. estimator_intro.ipynb pickling_cuml_models.ipynb dask_multigpu_guide.ipynb + advanced.rst health_checks.rst supported_versions.rst From a2d706e6c27964b3113eb58c0ba6d9704d0c6403 Mon Sep 17 00:00:00 2001 From: Philip Hyunsu Cho Date: Thu, 21 May 2026 14:10:34 -0700 Subject: [PATCH 16/17] Adopt nvForest for random forest inference (#8048) Adopts `nvForest` for cuML random forest inference, removes the in-tree FIL implementation, and keeps a deprecated `cuml.fil` compatibility layer that directs users to `nvforest.ForestInference`. This also wires `nvForest` into the C++ and Python build paths, updates packaging dependencies, and preserves the existing `cuml.fil.ForestInference` entry point with deprecation warnings for users migrating from FIL. Authors: - Philip Hyunsu Cho (https://github.com/hcho3) - Simon Adorf (https://github.com/csadorf) Approvers: - Simon Adorf (https://github.com/csadorf) - Jim Crist-Harif (https://github.com/jcrist) - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cuml/pull/8048 --- ci/build_wheel_cuml.sh | 7 +- ci/build_wheel_libcuml.sh | 5 +- ci/release/update-version.sh | 4 +- .../all_cuda-129_arch-aarch64.yaml | 2 + .../all_cuda-129_arch-x86_64.yaml | 2 + .../all_cuda-132_arch-aarch64.yaml | 2 + .../all_cuda-132_arch-x86_64.yaml | 2 + .../clang_tidy_cuda-129_arch-x86_64.yaml | 1 + .../clang_tidy_cuda-132_arch-x86_64.yaml | 1 + .../cpp_all_cuda-129_arch-x86_64.yaml | 1 + .../cpp_all_cuda-132_arch-x86_64.yaml | 1 + conda/recipes/cuml/recipe.yaml | 2 + conda/recipes/libcuml/recipe.yaml | 2 + cpp/CMakeLists.txt | 49 +- cpp/bench/CMakeLists.txt | 3 +- cpp/bench/sg/fil.cu | 203 --- cpp/cmake/modules/ConfigureAlgorithms.cmake | 4 +- cpp/cmake/thirdparty/get_nvforest.cmake | 52 + cpp/include/cuml/fil/Implementation.md | 232 --- cpp/include/cuml/fil/README.md | 161 -- cpp/include/cuml/fil/constants.hpp | 27 - cpp/include/cuml/fil/decision_forest.hpp | 483 ------ cpp/include/cuml/fil/detail/bitset.hpp | 110 -- .../cuml/fil/detail/cpu_introspection.hpp | 19 - .../fil/detail/decision_forest_builder.hpp | 361 ----- .../cuml/fil/detail/degenerate_trees.hpp | 75 - .../cuml/fil/detail/device_initialization.hpp | 36 - .../fil/detail/device_initialization/cpu.hpp | 34 - .../fil/detail/device_initialization/gpu.cuh | 242 --- .../fil/detail/device_initialization/gpu.hpp | 32 - cpp/include/cuml/fil/detail/evaluate_tree.hpp | 201 --- cpp/include/cuml/fil/detail/forest.hpp | 77 - .../cuml/fil/detail/gpu_introspection.hpp | 108 -- cpp/include/cuml/fil/detail/index_type.hpp | 11 - cpp/include/cuml/fil/detail/infer.hpp | 169 -- cpp/include/cuml/fil/detail/infer/cpu.hpp | 147 -- cpp/include/cuml/fil/detail/infer/gpu.cuh | 345 ----- cpp/include/cuml/fil/detail/infer/gpu.hpp | 45 - .../cuml/fil/detail/infer_kernel/cpu.hpp | 203 --- .../cuml/fil/detail/infer_kernel/gpu.cuh | 213 --- .../infer_kernel/shared_memory_buffer.cuh | 147 -- cpp/include/cuml/fil/detail/node.hpp | 248 --- cpp/include/cuml/fil/detail/postprocessor.hpp | 236 --- .../cuml/fil/detail/raft_proto/buffer.hpp | 390 ----- .../cuml/fil/detail/raft_proto/ceildiv.hpp | 17 - .../cuml/fil/detail/raft_proto/cuda_check.hpp | 19 - .../fil/detail/raft_proto/cuda_stream.hpp | 22 - .../raft_proto/detail/const_agnostic.hpp | 16 - .../fil/detail/raft_proto/detail/copy.hpp | 83 - .../fil/detail/raft_proto/detail/copy/cpu.hpp | 39 - .../fil/detail/raft_proto/detail/copy/gpu.hpp | 31 - .../raft_proto/detail/cuda_check/base.hpp | 17 - .../raft_proto/detail/cuda_check/gpu.hpp | 24 - .../raft_proto/detail/device_id/base.hpp | 18 - .../raft_proto/detail/device_id/cpu.hpp | 23 - .../raft_proto/detail/device_id/gpu.hpp | 31 - .../raft_proto/detail/device_setter/base.hpp | 19 - .../raft_proto/detail/device_setter/gpu.hpp | 38 - .../raft_proto/detail/host_only_throw.hpp | 13 - .../detail/host_only_throw/base.hpp | 19 - .../raft_proto/detail/host_only_throw/cpu.hpp | 20 - .../raft_proto/detail/non_owning_buffer.hpp | 12 - .../detail/non_owning_buffer/base.hpp | 28 - .../raft_proto/detail/owning_buffer.hpp | 14 - .../raft_proto/detail/owning_buffer/base.hpp | 23 - .../raft_proto/detail/owning_buffer/cpu.hpp | 31 - .../raft_proto/detail/owning_buffer/gpu.hpp | 41 - .../cuml/fil/detail/raft_proto/device_id.hpp | 21 - .../fil/detail/raft_proto/device_setter.hpp | 16 - .../fil/detail/raft_proto/device_type.hpp | 8 - .../cuml/fil/detail/raft_proto/exceptions.hpp | 56 - .../fil/detail/raft_proto/gpu_support.hpp | 45 - .../cuml/fil/detail/raft_proto/handle.hpp | 43 - .../cuml/fil/detail/raft_proto/padding.hpp | 48 - .../cuml/fil/detail/specialization_types.hpp | 76 - .../device_initialization_macros.hpp | 14 - .../detail/specializations/forest_macros.hpp | 27 - .../detail/specializations/infer_macros.hpp | 143 -- cpp/include/cuml/fil/exceptions.hpp | 64 - cpp/include/cuml/fil/forest_model.hpp | 306 ---- cpp/include/cuml/fil/infer_kind.hpp | 10 - cpp/include/cuml/fil/postproc_ops.hpp | 27 - cpp/include/cuml/fil/tree_layout.hpp | 19 - cpp/include/cuml/fil/treelite_importer.hpp | 509 ------ cpp/include/cuml/forest/README.md | 7 - cpp/include/cuml/forest/exceptions.hpp | 21 - .../cuml/forest/integrations/treelite.hpp | 226 --- .../forest/traversal/traversal_forest.hpp | 190 --- .../cuml/forest/traversal/traversal_node.hpp | 39 - .../cuml/forest/traversal/traversal_order.hpp | 38 - cpp/src/fil/infer0.cpp | 15 - cpp/src/fil/infer0.cu | 20 - cpp/src/fil/infer1.cpp | 15 - cpp/src/fil/infer1.cu | 20 - cpp/src/fil/infer10.cpp | 15 - cpp/src/fil/infer10.cu | 20 - cpp/src/fil/infer11.cpp | 15 - cpp/src/fil/infer11.cu | 20 - cpp/src/fil/infer2.cpp | 15 - cpp/src/fil/infer2.cu | 20 - cpp/src/fil/infer3.cpp | 15 - cpp/src/fil/infer3.cu | 20 - cpp/src/fil/infer4.cpp | 15 - cpp/src/fil/infer4.cu | 20 - cpp/src/fil/infer5.cpp | 15 - cpp/src/fil/infer5.cu | 20 - cpp/src/fil/infer6.cpp | 15 - cpp/src/fil/infer6.cu | 20 - cpp/src/fil/infer7.cpp | 15 - cpp/src/fil/infer7.cu | 20 - cpp/src/fil/infer8.cpp | 15 - cpp/src/fil/infer8.cu | 20 - cpp/src/fil/infer9.cpp | 15 - cpp/src/fil/infer9.cu | 20 - cpp/tests/CMakeLists.txt | 14 +- ...decision_forest_builder_invalid_inputs.cpp | 82 - cpp/tests/sg/fil/raft_proto/buffer.cpp | 379 ----- cpp/tests/sg/fil/raft_proto/buffer.cu | 42 - cpp/tests/sg/fil/treelite_importer.cpp | 378 ----- .../fil/treelite_importer_invalid_inputs.cpp | 248 --- cpp/tests/sg/forest/traversal_forest.cpp | 327 ---- cpp/tests/sg/forest/treelite_traversal.cpp | 487 ------ cpp/tests/sg/rf_test.cu | 137 +- dependencies.yaml | 64 + docs/source/FIL.rst | 235 +-- python/cuml/CMakeLists.txt | 1 - python/cuml/cuml/__init__.py | 3 +- python/cuml/cuml/benchmark/algorithms.py | 89 +- .../cuml/cuml/benchmark/bench_helper_funcs.py | 295 ---- .../cuml/ensemble/randomforest_common.pyx | 84 +- .../cuml/ensemble/randomforestclassifier.py | 32 +- .../cuml/ensemble/randomforestregressor.py | 20 +- python/cuml/cuml/fil/CMakeLists.txt | 19 - python/cuml/cuml/fil/__init__.py | 4 +- python/cuml/cuml/fil/compat.py | 514 ++++++ python/cuml/cuml/fil/detail/__init__.py | 4 - .../cuml/fil/detail/raft_proto/__init__.py | 4 - .../fil/detail/raft_proto/cuda_stream.pxd | 7 - .../fil/detail/raft_proto/device_type.pxd | 8 - .../cuml/fil/detail/raft_proto/handle.pxd | 19 - .../cuml/fil/detail/raft_proto/optional.pxd | 43 - python/cuml/cuml/fil/fil.pyx | 1371 ----------------- python/cuml/cuml/fil/infer_kind.pxd | 11 - python/cuml/cuml/fil/postprocessing.pxd | 16 - python/cuml/cuml/fil/tree_layout.pxd | 9 - python/cuml/pyproject.toml | 3 + .../tests/explainer/test_explainer_common.py | 14 +- python/cuml/tests/test_api.py | 5 + python/cuml/tests/test_benchmark.py | 25 +- python/cuml/tests/test_fil.py | 31 +- python/cuml/tests/test_random_forest.py | 41 +- python/libcuml/libcuml/load.py | 2 + python/libcuml/pyproject.toml | 2 + 153 files changed, 928 insertions(+), 11807 deletions(-) delete mode 100644 cpp/bench/sg/fil.cu create mode 100644 cpp/cmake/thirdparty/get_nvforest.cmake delete mode 100644 cpp/include/cuml/fil/Implementation.md delete mode 100644 cpp/include/cuml/fil/README.md delete mode 100644 cpp/include/cuml/fil/constants.hpp delete mode 100644 cpp/include/cuml/fil/decision_forest.hpp delete mode 100644 cpp/include/cuml/fil/detail/bitset.hpp delete mode 100644 cpp/include/cuml/fil/detail/cpu_introspection.hpp delete mode 100644 cpp/include/cuml/fil/detail/decision_forest_builder.hpp delete mode 100644 cpp/include/cuml/fil/detail/degenerate_trees.hpp delete mode 100644 cpp/include/cuml/fil/detail/device_initialization.hpp delete mode 100644 cpp/include/cuml/fil/detail/device_initialization/cpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/device_initialization/gpu.cuh delete mode 100644 cpp/include/cuml/fil/detail/device_initialization/gpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/evaluate_tree.hpp delete mode 100644 cpp/include/cuml/fil/detail/forest.hpp delete mode 100644 cpp/include/cuml/fil/detail/gpu_introspection.hpp delete mode 100644 cpp/include/cuml/fil/detail/index_type.hpp delete mode 100644 cpp/include/cuml/fil/detail/infer.hpp delete mode 100644 cpp/include/cuml/fil/detail/infer/cpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/infer/gpu.cuh delete mode 100644 cpp/include/cuml/fil/detail/infer/gpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/infer_kernel/cpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/infer_kernel/gpu.cuh delete mode 100644 cpp/include/cuml/fil/detail/infer_kernel/shared_memory_buffer.cuh delete mode 100644 cpp/include/cuml/fil/detail/node.hpp delete mode 100644 cpp/include/cuml/fil/detail/postprocessor.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/buffer.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/ceildiv.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/cuda_check.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/cuda_stream.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/const_agnostic.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/copy.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/copy/cpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/copy/gpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/cuda_check/base.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/cuda_check/gpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/device_id/base.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/device_id/cpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/device_id/gpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/device_setter/base.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/device_setter/gpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw/base.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw/cpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/non_owning_buffer.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/non_owning_buffer/base.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/base.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/cpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/gpu.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/device_id.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/device_setter.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/device_type.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/exceptions.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/gpu_support.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/handle.hpp delete mode 100644 cpp/include/cuml/fil/detail/raft_proto/padding.hpp delete mode 100644 cpp/include/cuml/fil/detail/specialization_types.hpp delete mode 100644 cpp/include/cuml/fil/detail/specializations/device_initialization_macros.hpp delete mode 100644 cpp/include/cuml/fil/detail/specializations/forest_macros.hpp delete mode 100644 cpp/include/cuml/fil/detail/specializations/infer_macros.hpp delete mode 100644 cpp/include/cuml/fil/exceptions.hpp delete mode 100644 cpp/include/cuml/fil/forest_model.hpp delete mode 100644 cpp/include/cuml/fil/infer_kind.hpp delete mode 100644 cpp/include/cuml/fil/postproc_ops.hpp delete mode 100644 cpp/include/cuml/fil/tree_layout.hpp delete mode 100644 cpp/include/cuml/fil/treelite_importer.hpp delete mode 100644 cpp/include/cuml/forest/README.md delete mode 100644 cpp/include/cuml/forest/exceptions.hpp delete mode 100644 cpp/include/cuml/forest/integrations/treelite.hpp delete mode 100644 cpp/include/cuml/forest/traversal/traversal_forest.hpp delete mode 100644 cpp/include/cuml/forest/traversal/traversal_node.hpp delete mode 100644 cpp/include/cuml/forest/traversal/traversal_order.hpp delete mode 100644 cpp/src/fil/infer0.cpp delete mode 100644 cpp/src/fil/infer0.cu delete mode 100644 cpp/src/fil/infer1.cpp delete mode 100644 cpp/src/fil/infer1.cu delete mode 100644 cpp/src/fil/infer10.cpp delete mode 100644 cpp/src/fil/infer10.cu delete mode 100644 cpp/src/fil/infer11.cpp delete mode 100644 cpp/src/fil/infer11.cu delete mode 100644 cpp/src/fil/infer2.cpp delete mode 100644 cpp/src/fil/infer2.cu delete mode 100644 cpp/src/fil/infer3.cpp delete mode 100644 cpp/src/fil/infer3.cu delete mode 100644 cpp/src/fil/infer4.cpp delete mode 100644 cpp/src/fil/infer4.cu delete mode 100644 cpp/src/fil/infer5.cpp delete mode 100644 cpp/src/fil/infer5.cu delete mode 100644 cpp/src/fil/infer6.cpp delete mode 100644 cpp/src/fil/infer6.cu delete mode 100644 cpp/src/fil/infer7.cpp delete mode 100644 cpp/src/fil/infer7.cu delete mode 100644 cpp/src/fil/infer8.cpp delete mode 100644 cpp/src/fil/infer8.cu delete mode 100644 cpp/src/fil/infer9.cpp delete mode 100644 cpp/src/fil/infer9.cu delete mode 100644 cpp/tests/sg/fil/decision_forest_builder_invalid_inputs.cpp delete mode 100644 cpp/tests/sg/fil/raft_proto/buffer.cpp delete mode 100644 cpp/tests/sg/fil/raft_proto/buffer.cu delete mode 100644 cpp/tests/sg/fil/treelite_importer.cpp delete mode 100644 cpp/tests/sg/fil/treelite_importer_invalid_inputs.cpp delete mode 100644 cpp/tests/sg/forest/traversal_forest.cpp delete mode 100644 cpp/tests/sg/forest/treelite_traversal.cpp delete mode 100644 python/cuml/cuml/fil/CMakeLists.txt create mode 100644 python/cuml/cuml/fil/compat.py delete mode 100644 python/cuml/cuml/fil/detail/__init__.py delete mode 100644 python/cuml/cuml/fil/detail/raft_proto/__init__.py delete mode 100644 python/cuml/cuml/fil/detail/raft_proto/cuda_stream.pxd delete mode 100644 python/cuml/cuml/fil/detail/raft_proto/device_type.pxd delete mode 100644 python/cuml/cuml/fil/detail/raft_proto/handle.pxd delete mode 100644 python/cuml/cuml/fil/detail/raft_proto/optional.pxd delete mode 100644 python/cuml/cuml/fil/fil.pyx delete mode 100644 python/cuml/cuml/fil/infer_kind.pxd delete mode 100644 python/cuml/cuml/fil/postprocessing.pxd delete mode 100644 python/cuml/cuml/fil/tree_layout.pxd diff --git a/ci/build_wheel_cuml.sh b/ci/build_wheel_cuml.sh index 422bd98525..d1f34c4c4c 100755 --- a/ci/build_wheel_cuml.sh +++ b/ci/build_wheel_cuml.sh @@ -20,16 +20,17 @@ LIBCUML_WHEELHOUSE=$(RAPIDS_PY_WHEEL_NAME="libcuml_${RAPIDS_PY_CUDA_SUFFIX}" rap echo "libcuml-${RAPIDS_PY_CUDA_SUFFIX} @ file://$(echo "${LIBCUML_WHEELHOUSE}"/libcuml_*.whl)" >> "${PIP_CONSTRAINT}" EXCLUDE_ARGS=( - --exclude "libcuml.so" - --exclude "libcuvs.so" - --exclude "libraft.so" --exclude "libcublas.so.*" --exclude "libcublasLt.so.*" --exclude "libcufft.so.*" + --exclude "libcuml.so" --exclude "libcurand.so.*" --exclude "libcusolver.so.*" --exclude "libcusparse.so.*" + --exclude "libcuvs.so" + --exclude "libnvforest++.so" --exclude "libnvJitLink.so.*" + --exclude "libraft.so" --exclude "librapids_logger.so" --exclude "librmm.so" ) diff --git a/ci/build_wheel_libcuml.sh b/ci/build_wheel_libcuml.sh index 2e2f2e35bf..2a31b40cca 100755 --- a/ci/build_wheel_libcuml.sh +++ b/ci/build_wheel_libcuml.sh @@ -29,17 +29,18 @@ rapids-pip-retry install \ export PIP_NO_BUILD_ISOLATION=0 EXCLUDE_ARGS=( - --exclude "libraft.so" --exclude "libcublas.so.*" --exclude "libcublasLt.so.*" --exclude "libcufft.so.*" --exclude "libcurand.so.*" --exclude "libcusolver.so.*" --exclude "libcusparse.so.*" + --exclude "libnccl.so.*" + --exclude "libnvforest++.so" --exclude "libnvJitLink.so.*" + --exclude "libraft.so" --exclude "librapids_logger.so" --exclude "librmm.so" - --exclude "libnccl.so.*" ) export SKBUILD_CMAKE_ARGS="-DDISABLE_DEPRECATION_WARNINGS=ON;-DCUML_USE_CUVS_STATIC=ON" diff --git a/ci/release/update-version.sh b/ci/release/update-version.sh index 081c9f2e6e..4d161b25df 100755 --- a/ci/release/update-version.sh +++ b/ci/release/update-version.sh @@ -1,5 +1,5 @@ #!/bin/bash -# SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 ######################## # cuML Version Updater # @@ -106,9 +106,11 @@ DEPENDENCIES=( libcuml libcuml-tests libcuvs + libnvforest libraft-headers libraft librmm + nvforest pylibraft raft-dask rapids-dask-dependency diff --git a/conda/environments/all_cuda-129_arch-aarch64.yaml b/conda/environments/all_cuda-129_arch-aarch64.yaml index f70b17a786..d636da6bc6 100644 --- a/conda/environments/all_cuda-129_arch-aarch64.yaml +++ b/conda/environments/all_cuda-129_arch-aarch64.yaml @@ -37,6 +37,7 @@ dependencies: - libcusolver-dev - libcusparse-dev - libcuvs==26.6.*,>=0.0.0a0 +- libnvforest==26.6.*,>=0.0.0a0 - libopenblas<=0.3.30 - libraft==26.6.*,>=0.0.0a0 - librmm==26.6.*,>=0.0.0a0 @@ -50,6 +51,7 @@ dependencies: - numpy>=1.26,<3.0 - numpydoc - numpydoc<1.9 +- nvforest==26.6.*,>=0.0.0a0 - nvidia-ml-py>=12 - onnxruntime - packaging diff --git a/conda/environments/all_cuda-129_arch-x86_64.yaml b/conda/environments/all_cuda-129_arch-x86_64.yaml index 130a281624..04dc51f95f 100644 --- a/conda/environments/all_cuda-129_arch-x86_64.yaml +++ b/conda/environments/all_cuda-129_arch-x86_64.yaml @@ -37,6 +37,7 @@ dependencies: - libcusolver-dev - libcusparse-dev - libcuvs==26.6.*,>=0.0.0a0 +- libnvforest==26.6.*,>=0.0.0a0 - libraft==26.6.*,>=0.0.0a0 - librmm==26.6.*,>=0.0.0a0 - lightgbm @@ -49,6 +50,7 @@ dependencies: - numpy>=1.26,<3.0 - numpydoc - numpydoc<1.9 +- nvforest==26.6.*,>=0.0.0a0 - nvidia-ml-py>=12 - onnxruntime - packaging diff --git a/conda/environments/all_cuda-132_arch-aarch64.yaml b/conda/environments/all_cuda-132_arch-aarch64.yaml index 40981932cf..0a674b9e97 100644 --- a/conda/environments/all_cuda-132_arch-aarch64.yaml +++ b/conda/environments/all_cuda-132_arch-aarch64.yaml @@ -37,6 +37,7 @@ dependencies: - libcusolver-dev - libcusparse-dev - libcuvs==26.6.*,>=0.0.0a0 +- libnvforest==26.6.*,>=0.0.0a0 - libopenblas<=0.3.30 - libraft==26.6.*,>=0.0.0a0 - librmm==26.6.*,>=0.0.0a0 @@ -50,6 +51,7 @@ dependencies: - numpy>=1.26,<3.0 - numpydoc - numpydoc<1.9 +- nvforest==26.6.*,>=0.0.0a0 - nvidia-ml-py>=12 - onnxruntime - packaging diff --git a/conda/environments/all_cuda-132_arch-x86_64.yaml b/conda/environments/all_cuda-132_arch-x86_64.yaml index b879ce0d48..b861ffd179 100644 --- a/conda/environments/all_cuda-132_arch-x86_64.yaml +++ b/conda/environments/all_cuda-132_arch-x86_64.yaml @@ -37,6 +37,7 @@ dependencies: - libcusolver-dev - libcusparse-dev - libcuvs==26.6.*,>=0.0.0a0 +- libnvforest==26.6.*,>=0.0.0a0 - libraft==26.6.*,>=0.0.0a0 - librmm==26.6.*,>=0.0.0a0 - lightgbm @@ -49,6 +50,7 @@ dependencies: - numpy>=1.26,<3.0 - numpydoc - numpydoc<1.9 +- nvforest==26.6.*,>=0.0.0a0 - nvidia-ml-py>=12 - onnxruntime - packaging diff --git a/conda/environments/clang_tidy_cuda-129_arch-x86_64.yaml b/conda/environments/clang_tidy_cuda-129_arch-x86_64.yaml index 10ddcdafa3..6860700446 100644 --- a/conda/environments/clang_tidy_cuda-129_arch-x86_64.yaml +++ b/conda/environments/clang_tidy_cuda-129_arch-x86_64.yaml @@ -21,6 +21,7 @@ dependencies: - libcusolver-dev - libcusparse-dev - libcuvs==26.6.*,>=0.0.0a0 +- libnvforest==26.6.*,>=0.0.0a0 - libraft-headers==26.6.*,>=0.0.0a0 - librmm==26.6.*,>=0.0.0a0 - llvm-openmp==20.1.8 diff --git a/conda/environments/clang_tidy_cuda-132_arch-x86_64.yaml b/conda/environments/clang_tidy_cuda-132_arch-x86_64.yaml index 5925bc6a3e..ec30c70ab7 100644 --- a/conda/environments/clang_tidy_cuda-132_arch-x86_64.yaml +++ b/conda/environments/clang_tidy_cuda-132_arch-x86_64.yaml @@ -21,6 +21,7 @@ dependencies: - libcusolver-dev - libcusparse-dev - libcuvs==26.6.*,>=0.0.0a0 +- libnvforest==26.6.*,>=0.0.0a0 - libraft-headers==26.6.*,>=0.0.0a0 - librmm==26.6.*,>=0.0.0a0 - llvm-openmp==20.1.8 diff --git a/conda/environments/cpp_all_cuda-129_arch-x86_64.yaml b/conda/environments/cpp_all_cuda-129_arch-x86_64.yaml index dd26bfc1c5..ef0fc258e6 100644 --- a/conda/environments/cpp_all_cuda-129_arch-x86_64.yaml +++ b/conda/environments/cpp_all_cuda-129_arch-x86_64.yaml @@ -19,6 +19,7 @@ dependencies: - libcusolver-dev - libcusparse-dev - libcuvs==26.6.*,>=0.0.0a0 +- libnvforest==26.6.*,>=0.0.0a0 - libraft-headers==26.6.*,>=0.0.0a0 - librmm==26.6.*,>=0.0.0a0 - ninja diff --git a/conda/environments/cpp_all_cuda-132_arch-x86_64.yaml b/conda/environments/cpp_all_cuda-132_arch-x86_64.yaml index d10c3cec61..0768d40fc4 100644 --- a/conda/environments/cpp_all_cuda-132_arch-x86_64.yaml +++ b/conda/environments/cpp_all_cuda-132_arch-x86_64.yaml @@ -19,6 +19,7 @@ dependencies: - libcusolver-dev - libcusparse-dev - libcuvs==26.6.*,>=0.0.0a0 +- libnvforest==26.6.*,>=0.0.0a0 - libraft-headers==26.6.*,>=0.0.0a0 - librmm==26.6.*,>=0.0.0a0 - ninja diff --git a/conda/recipes/cuml/recipe.yaml b/conda/recipes/cuml/recipe.yaml index 960d4347df..710cde8e05 100644 --- a/conda/recipes/cuml/recipe.yaml +++ b/conda/recipes/cuml/recipe.yaml @@ -80,6 +80,7 @@ requirements: - cudf =${{ minor_version }} - cython >=3.2.2 - libcuml =${{ version }} + - nvforest =${{ minor_version }} - pip - pylibraft =${{ minor_version }} - python =${{ py_abi_min }} @@ -101,6 +102,7 @@ requirements: - numba >=0.60.0,<0.65.0 - numba-cuda >=0.22.2,<0.29.0 - numpy >=1.26,<3.0 + - nvforest =${{ minor_version }} - scikit-learn >=1.4 - scipy >=1.14.0 - packaging diff --git a/conda/recipes/libcuml/recipe.yaml b/conda/recipes/libcuml/recipe.yaml index 5d9348a7b6..cd0f61b750 100644 --- a/conda/recipes/libcuml/recipe.yaml +++ b/conda/recipes/libcuml/recipe.yaml @@ -106,6 +106,7 @@ outputs: - cuda-version =${{ cuda_version }} - rapids-logger =0.2 - libcuvs =${{ minor_version }} + - libnvforest =${{ minor_version }} - librmm =${{ minor_version }} - treelite ${{ treelite_version }} - cuda-cudart-dev @@ -122,6 +123,7 @@ outputs: - libcusolver - libcusparse - libcuvs =${{ minor_version }} + - libnvforest =${{ minor_version }} - librmm =${{ minor_version }} - rapids-logger =0.2 - treelite ${{ treelite_version }} diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 617a45af18..f05dfad14b 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -70,8 +70,12 @@ option(CUML_COMPILE_DYNAMIC_ONLY "Only build the shared library and skip the sta # dependent library. option(CUML_EXCLUDE_RAFT_FROM_ALL "Exclude RAFT targets from cuML's 'all' target" OFF) option(CUML_EXCLUDE_TREELITE_FROM_ALL "Exclude Treelite targets from cuML's 'all' target" OFF) +option(CUML_EXCLUDE_NVFOREST_FROM_ALL "Exclude nvForest targets from cuML's 'all' target" OFF) option(CUML_RAFT_CLONE_ON_PIN "Explicitly clone RAFT branch when pinned to non-feature branch" ON) option(CUML_CUVS_CLONE_ON_PIN "Explicitly clone CUVS branch when pinned to non-feature branch" ON) +option(CUML_NVFOREST_CLONE_ON_PIN + "Explicitly clone nvForest branch when pinned to non-feature branch" ON +) message(VERBOSE "CUML_CPP: Building libcuml shared library: ${BUILD_CUML_CPP_LIBRARY}") message(VERBOSE "CUML_CPP: Building cuML algorithm tests: ${BUILD_CUML_TESTS}") @@ -270,6 +274,11 @@ if(LINK_TREELITE) include(cmake/thirdparty/get_treelite.cmake) endif() +if(LINK_NVFOREST) + include(cmake/thirdparty/get_nvforest.cmake) + set(CUML_NVFOREST_TARGET nvforest::nvforest++) +endif() + if(all_algo OR treeshap_algo) include(cmake/thirdparty/get_gputreeshap.cmake) endif() @@ -367,42 +376,6 @@ if(BUILD_CUML_CPP_LIBRARY) target_sources(cuml_objs PRIVATE src/explainer/tree_shap.cu) endif() - # FIL components - if(all_algo OR fil_algo) - if(CUML_ENABLE_GPU) - target_sources( - cuml_objs - PRIVATE src/fil/infer0.cu - src/fil/infer1.cu - src/fil/infer2.cu - src/fil/infer3.cu - src/fil/infer4.cu - src/fil/infer5.cu - src/fil/infer6.cu - src/fil/infer7.cu - src/fil/infer8.cu - src/fil/infer9.cu - src/fil/infer10.cu - src/fil/infer11.cu - ) - endif() - target_sources( - cuml_objs - PRIVATE src/fil/infer0.cpp - src/fil/infer1.cpp - src/fil/infer2.cpp - src/fil/infer3.cpp - src/fil/infer4.cpp - src/fil/infer5.cpp - src/fil/infer6.cpp - src/fil/infer7.cpp - src/fil/infer8.cpp - src/fil/infer9.cpp - src/fil/infer10.cpp - src/fil/infer11.cpp - ) - endif() - # todo: organize linear models better if(all_algo OR linearregression_algo @@ -596,6 +569,10 @@ if(BUILD_CUML_CPP_LIBRARY) copy_interface_excludes(INCLUDED_TARGET treelite::treelite_static TARGET cuml_objs) endif() + if(LINK_NVFOREST) + list(APPEND _cuml_cpp_public_libs ${CUML_NVFOREST_TARGET}) + endif() + # These are always private: list( APPEND diff --git a/cpp/bench/CMakeLists.txt b/cpp/bench/CMakeLists.txt index 727dc15cad..24d576237f 100644 --- a/cpp/bench/CMakeLists.txt +++ b/cpp/bench/CMakeLists.txt @@ -1,6 +1,6 @@ # ============================================================================= # cmake-format: off -# SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # cmake-format: on # ============================================================================= @@ -24,7 +24,6 @@ if(BUILD_CUML_BENCH) sg/svc.cu sg/svr.cu sg/umap.cu - sg/fil.cu ) if(CUML_ENABLE_GPU) target_compile_definitions(${CUML_CPP_BENCH_TARGET} PUBLIC CUML_ENABLE_GPU) diff --git a/cpp/bench/sg/fil.cu b/cpp/bench/sg/fil.cu deleted file mode 100644 index 129a2ff6ec..0000000000 --- a/cpp/bench/sg/fil.cu +++ /dev/null @@ -1,203 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include "benchmark.cuh" - -#include -#include -#include -#include -#include -#include -#include - -#include - -#include -#include -#include - -namespace ML { -namespace Bench { -namespace fil { - -struct Params { - DatasetParams data; - RegressionParams blobs; - TreeliteModelHandle model; - RF_params rf; - int predict_repetitions; -}; - -class FIL : public RegressionFixture { - typedef RegressionFixture Base; - - public: - FIL(const std::string& name, const Params& p) - : RegressionFixture(name, p.data, p.blobs), model(p.model), p_rest(p) - { - } - - protected: - void runBenchmark(::benchmark::State& state) override - { - if (!params.rowMajor) { state.SkipWithError("FIL only supports row-major inputs"); } - // create model - ML::RandomForestRegressorF rf_model; - auto* mPtr = &rf_model; - auto train_nrows = std::min(params.nrows, 1000); - fit(*handle, mPtr, data.X.data(), train_nrows, params.ncols, data.y.data(), p_rest.rf); - handle->sync_stream(stream); - - ML::build_treelite_forest(&model, &rf_model, params.ncols); - - auto fil_model = ML::fil::import_from_treelite_handle(model, - ML::fil::tree_layout::breadth_first, - 128, - false, - raft_proto::device_type::gpu, - 0, - stream); - - auto optimal_chunk_size = 1; - auto optimal_layout = ML::fil::tree_layout::breadth_first; - auto allowed_layouts = - std::vector{ML::fil::tree_layout::depth_first, - ML::fil::tree_layout::breadth_first, - ML::fil::tree_layout::layered_children_together}; - auto min_time = std::numeric_limits::max(); - - // Find optimal configuration - for (auto layout : allowed_layouts) { - fil_model = ML::fil::import_from_treelite_handle( - model, layout, 128, false, raft_proto::device_type::gpu, 0, stream); - for (auto chunk_size = 1; chunk_size <= 32; chunk_size *= 2) { - handle->sync_stream(); - handle->sync_stream_pool(); - auto start = std::chrono::high_resolution_clock::now(); - for (int i = 0; i < p_rest.predict_repetitions; i++) { - // Create FIL forest - fil_model.predict(*handle, - data.y.data(), - data.X.data(), - params.nrows, - raft_proto::device_type::gpu, - raft_proto::device_type::gpu, - ML::fil::infer_kind::default_kind, - chunk_size); - } - handle->sync_stream(); - handle->sync_stream_pool(); - auto end = std::chrono::high_resolution_clock::now(); - auto elapsed = std::chrono::duration_cast(end - start).count(); - if (elapsed < min_time) { - min_time = elapsed; - optimal_chunk_size = chunk_size; - optimal_layout = layout; - } - } - } - - // Build optimal FIL tree - fil_model = ML::fil::import_from_treelite_handle( - model, optimal_layout, 128, false, raft_proto::device_type::gpu, 0, stream); - - handle->sync_stream(); - handle->sync_stream_pool(); - - // only time prediction - this->loopOnState( - state, - [this, &fil_model, optimal_chunk_size]() { - for (int i = 0; i < p_rest.predict_repetitions; i++) { - fil_model.predict(*handle, - this->data.y.data(), - this->data.X.data(), - this->params.nrows, - raft_proto::device_type::gpu, - raft_proto::device_type::gpu, - ML::fil::infer_kind::default_kind, - optimal_chunk_size); - handle->sync_stream(); - handle->sync_stream_pool(); - } - }, - true); - } - - void allocateBuffers(const ::benchmark::State& state) override { Base::allocateBuffers(state); } - - void deallocateBuffers(const ::benchmark::State& state) override - { - Base::deallocateBuffers(state); - } - - private: - TreeliteModelHandle model; - Params p_rest; -}; - -struct FilBenchParams { - int nrows; - int ncols; - int nclasses; - int max_depth; - int ntrees; -}; - -std::vector getInputs() -{ - std::vector out; - Params p; - p.data.rowMajor = true; - p.blobs = {.n_informative = -1, // Just a placeholder value, anyway changed below - .effective_rank = -1, // Just a placeholder value, anyway changed below - .bias = 0.f, - .tail_strength = 0.1, - .noise = 0.01, - .shuffle = false, - .seed = 12345ULL}; - - p.rf = set_rf_params(10, /*max_depth */ - (1 << 20), /* max_leaves */ - 1.f, /* max_features */ - 32, /* max_n_bins */ - 3, /* min_samples_leaf */ - 3, /* min_samples_split */ - 0.0f, /* min_impurity_decrease */ - true, /* bootstrap */ - 1, /* n_trees */ - 1.f, /* max_samples */ - 1234ULL, /* seed */ - ML::CRITERION::MSE, /* split_criterion */ - 8, /* n_streams */ - 128 /* max_batch_size */ - ); - - std::vector var_params = {{(int)1e6, 20, 1, 10, 1000}, - {(int)1e6, 20, 1, 3, 1000}, - {(int)1e6, 20, 1, 28, 1000}, - {(int)1e6, 20, 1, 10, 100}, - {(int)1e6, 20, 1, 10, 10000}, - {(int)1e6, 200, 1, 10, 1000}}; - for (auto& i : var_params) { - p.data.nrows = i.nrows; - p.data.ncols = i.ncols; - p.blobs.n_informative = i.ncols / 3; - p.blobs.effective_rank = i.ncols / 3; - p.data.nclasses = i.nclasses; - p.rf.tree_params.max_depth = i.max_depth; - p.rf.n_trees = i.ntrees; - p.predict_repetitions = 10; - out.push_back(p); - } - return out; -} - -ML_BENCH_REGISTER(Params, FIL, "", getInputs()); - -} // namespace fil -} // end namespace Bench -} // end namespace ML diff --git a/cpp/cmake/modules/ConfigureAlgorithms.cmake b/cpp/cmake/modules/ConfigureAlgorithms.cmake index e5338bae7e..66494f52ad 100644 --- a/cpp/cmake/modules/ConfigureAlgorithms.cmake +++ b/cpp/cmake/modules/ConfigureAlgorithms.cmake @@ -10,6 +10,7 @@ if(CUML_ALGORITHMS STREQUAL "ALL") set(LINK_TREELITE ON) set(LINK_CUFFT ON) set(LINK_CUVS ON) + set(LINK_NVFOREST ON) set(all_algo ON) # setting treeshap to ON to get the gputreeshap include in the cuml target set(treeshap_algo ON) @@ -76,7 +77,7 @@ else() # Set linking options and algorithms that require other algorithms ####### - if(fil_algo OR treeshap_algo) + if(treeshap_algo) set(LINK_TREELITE ON) endif() @@ -102,6 +103,7 @@ else() if(randomforest_algo) set(decisiontree_algo ON) set(LINK_TREELITE ON) + set(LINK_NVFOREST ON) endif() if(hierarchicalclustering_algo OR kmeans_algo) diff --git a/cpp/cmake/thirdparty/get_nvforest.cmake b/cpp/cmake/thirdparty/get_nvforest.cmake new file mode 100644 index 0000000000..b2129aae96 --- /dev/null +++ b/cpp/cmake/thirdparty/get_nvforest.cmake @@ -0,0 +1,52 @@ +#============================================================================= +# cmake-format: off +# SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. +# SPDX-License-Identifier: Apache-2.0 +# cmake-format: on +#============================================================================= + +set(CUML_MIN_VERSION_nvforest "${CUML_VERSION_MAJOR}.${CUML_VERSION_MINOR}.00") + +function(find_and_configure_nvforest) + set(oneValueArgs VERSION FORK PINNED_TAG EXCLUDE_FROM_ALL CLONE_ON_PIN) + cmake_parse_arguments(PKG "${options}" "${oneValueArgs}" + "${multiValueArgs}" ${ARGN} ) + + if(PKG_CLONE_ON_PIN AND NOT PKG_PINNED_TAG STREQUAL "${rapids-cmake-checkout-tag}") + message(STATUS "CUML: nvForest pinned tag found: ${PKG_PINNED_TAG}. Cloning nvForest locally.") + set(CPM_DOWNLOAD_nvforest ON) + endif() + + rapids_cpm_find(nvforest ${PKG_VERSION} + GLOBAL_TARGETS nvforest::nvforest++ + BUILD_EXPORT_SET cuml-exports + INSTALL_EXPORT_SET cuml-exports + CPM_ARGS + GIT_REPOSITORY https://github.com/${PKG_FORK}/nvforest.git + GIT_TAG ${PKG_PINNED_TAG} + SOURCE_SUBDIR cpp + EXCLUDE_FROM_ALL ${PKG_EXCLUDE_FROM_ALL} + OPTIONS + "BUILD_NVFOREST_TESTS OFF" + ) + + if(nvforest_ADDED) + message(VERBOSE "CUML: Using nvForest located in ${nvforest_SOURCE_DIR}") + else() + message(VERBOSE "CUML: Using nvForest located in ${nvforest_DIR}") + endif() + +endfunction() + +# Change pinned tag here to test a commit in CI +# To use a different nvForest locally, set the CMake variable +# CPM_nvforest_SOURCE=/path/to/local/nvforest +find_and_configure_nvforest(VERSION ${CUML_MIN_VERSION_nvforest} + FORK rapidsai + PINNED_TAG ${rapids-cmake-checkout-tag} + EXCLUDE_FROM_ALL ${CUML_EXCLUDE_NVFOREST_FROM_ALL} + # When PINNED_TAG above doesn't match cuml, + # force local nvforest clone in build directory + # even if it's already installed. + CLONE_ON_PIN ${CUML_NVFOREST_CLONE_ON_PIN} +) diff --git a/cpp/include/cuml/fil/Implementation.md b/cpp/include/cuml/fil/Implementation.md deleted file mode 100644 index a6addd6c23..0000000000 --- a/cpp/include/cuml/fil/Implementation.md +++ /dev/null @@ -1,232 +0,0 @@ -# FIL Implementation -This document is intended to provide additional detail about this -implementation of FIL to help guide future FIL contributors. Because this is -the first cuML algorithm to attempt to provide a unified CPU/GPU codebase that -does *not* require nvcc, CUDA or any other GPU-related library for its CPU-only -build, we also go over general strategies for CPU/GPU interoperability as used -by FIL. - -**A NOTE ON THE `raft_proto` NAMESPACE:** In addition to FIL-specific code, the new -implementation requires some more general-purpose CPU-GPU interoperable -utilities. Many of these utilities are either already implemented in RAFT (but -do not provide the required CPU-interoperable compilation guarantees) or are a -natural fit for incorporation in RAFT. In order to allow for more careful -integration with the existing RAFT codebase and interoperability -strategies, these utilities are currently provided in the `raft_proto` -namespace but will be moved into RAFT over time. Other algorithms should -not make use of the `raft_proto` namespace but instead wait until this -transition has taken place. - -## Design Goals -1. Provide state-of-the-art runtime performance for forest models on GPU, - especially for cases where CPU performance will not suffice (e.g. large - batches, deep trees, many trees, etc.). -2. Ensure that the public API is the same for both CPU and GPU execution. -3. Re-use as much code as possible between CPU and GPU implementations. -4. Provide near-state-of-the-art runtime performance for forest models on most - CPUs without vendor-specific optimizations. - -## Strategies for CPU/GPU code re-use - -This FIL implementation now makes use of a build-time variable -`CUML_ENABLE_GPU` to determine whether or not to compile CUDA code. If -`CUML_ENABLE_GPU` is not set, FIL is guaranteed to compile without nvcc, access -to CUDA headers, or any other GPU-related library. - -We explicitly wish to avoid excessive use of `#ifdef` statements based on this -variable, however. Interleaving CPU and GPU code via `#ifdef` branches both -reduces readability and discourages writing of truly interoperable code. -Ideally, `#ifdef` statements should be used solely and sparingly for -conditional header inclusion. This presents additional challenges but also -opportunities for a cleaner implementation of a unified CPU/GPU -codebase. - -It is also occasionally useful to make use of a `constexpr` value -indicating whether or not `CUML_ENABLE_GPU` is set, which we introduce as -`raft_proto::GPU_ENABLED`. - -### Avoiding CUDA symbols in CPU-only builds -The most significant challenge of attempting to create a unified CPU/GPU -implementation is ensuring that no CUDA symbols are exposed in the CPU-only -build. To illustrate the general strategy, we will look at a specific example: -the implementation of the main inference loop. Code for this loop is provided -in the following four files: - -``` -detail/ -├─ infer.hpp # "Consumable" header -├─ infer/ # "Implementation" directory -│ ├─ cpu.hpp -│ ├─ gpu.cuh -│ ├─ gpu.hpp -``` - -For brevity, we introduce the concepts of "consumable" and "implementation" -headers. Consumable headers can be included in any other header and are -guaranteed not to themselves include any header with CUDA symbols -if `CUML_ENABLE_GPU` is not defined. -Implementation headers can *only* be included by their associated consumable -header or directly in a source file. They should *never* be directly included -by any other consumable header except their own. - -By creating a clear separation of these two header types, we guarantee that any -source file that includes a consumable header should be compilable with or -without access to CUDA headers. Note that all public headers should be -consumable, but not all consumable headers need be made public. In the -particular example under consideration, `infer.hpp` is consumable, but we keep -it in the detail directory to indicate that it is not part of the public API. - -Let's take a closer look at each of the "infer" headers. `infer.hpp` -implements `detail::infer`, a function templated on both the execution device -type (`D`) and the type of the forest model being evaluated `forest_t`. -If we were to look at the implementation of this template, we would note -that there is no code specialized for either possible value of `D`. At the -level of consumable headers, we have abstracted away the difference between -GPU and CPU in order to ensure that this template is completely reusable -between GPU and CPU. - -Where we _need_ to provide distinct logic between GPU and CPU -implementations, we do so in implementation headers. In `infer/cpu.hpp`, we -have a fully-defined template for CPU specializations of -`detail::inference::infer`. If `raft_proto::GPU_ENABLED` is `false`, we also -include the GPU specializations, which will simply throw an exception if -invoked. In `infer/gpu.hpp` we *declare* but do not *define* the GPU -specializations. In `infer/gpu.cuh` we provide the full working definition for -the GPU specializations. - -`infer.hpp` includes `infer/cpu.hpp` and `infer/gpu.hpp`, but *not* -`infer/gpu.cuh`. Instead, `infer/gpu.cuh` is included directly in the CUDA -source files that require access to the full definition. - -Structuring the code in this way, we have a single separation point between -code that will and will not compile without access to CUDA headers. A similar -approach is used anywhere else in the implementation where we need distinct -logic for CPU and GPU. Otherwise, we are free to use anything defined in a -consumable header without worrying about whether the current translation unit -will ultimately be compiled with gcc or nvcc or whether our current build does -or does not have GPU enabled. - -### Re-using code - -Ultimately, many GPU and parallel CPU algorithms do not differ much in their -actual steps, but optimizing each requires careful attention to the -differing parallelism models and memory access models on each hardware -type. This means that with a little care, we can separate details related -to parallelism and memory access from the actual algorithm logic. This logic -will be the same for both CPU and GPU, but the now-isolated parallelism -and memory access code can be independently optimized. - -The process of actually performing this separation usually starts by -identifying the basic single "task" that each parallel worker must take on. -It is not always entirely obvious how granular this task should be. For -instance, in the case of forest models, we might consider the basic task to -be evaluating a single row with all trees in the forest, evaluating all rows -with a single tree of the forest, evaluating a single row with a single -tree, evaluating a single node of a tree on a single row, evaluating a sub-tree -of a specific size on a single row, etc. - -In order to offer optimal performance on the widest range of models, the -present implementation defines the underlying worker task as evaluating -a single row on a single tree, but specific model characteristics (e.g. very -small or large trees) might benefit from other task granularity. - -Once we have identified the underlying task, we implement this -directly in a way that is independent of the parallelism model or memory -access patterns. That is to say, we assume that we are already executing -on a single worker and that the memory is arranged optimally for this task. In -the current implementation, this task is defined in -`detail/evaluate_tree.hpp`. - -Looking at this header, we should note that there is no logic specific to the -GPU or CPU. Instead we defer this to `infer_kernel`, which specifies how our -fundamental task gets assigned to individual "workers" (CPU threads for the CPU -or CUDA threads for the GPU). This is not a necessary constraint (i.e. we could -refactor later for CPU and GPU specific versions of `evaluate_tree`), but -re-using code in this way and providing a clean separation from the parallelism -model does offer advantages. - -Beyond just the reduced maintenance of a single codebase and more modular -design, this gives us the opportunity to benefit from improvements in the CPU -implementation on GPU and vice versa. For instance, during the initial -development, only CPU tests were used to check for correctness, but GPU results -were shown to be correct as soon as they were added to the tests. Similarly, -during optimization, only GPU runtime and instructions were analyzed, but the -process of optimizing for the GPU resulted in significant speedups (over 50% on -a standard benchmark) on the CPU. - -## Code Walkthrough - -With some motivation for the general approach to CPU-GPU interoperability, we -now offer an overview of the layout of the codebase to help guide future -improvements. Because `raft_proto` utilities are going to be moved to RAFT or other -general-purpose libraries, we will not review anything within the `raft_proto` -directory here. - -### Public Headers -* `constants.hpp`: Contains constant values that may be useful in working - with FIL in other C++ applications -* `decision_forest.hpp`: Provides `decision_forest`, a template which provides - concrete implementations of a decision forest. Because different types may - be optimal for different sizes of models or models with different features, - we implement this template on many different combinations of template - parameters. This is provided in a public header in case other - applications have more specialized use cases and can afford to work directly - with this concrete underlying object. -* `forest_model.hpp`: Provides `forest_model`, a wrapper for a `std::variant` - of all `decision_forest` implementations. This wrapper handles - dispatching `predict` calls to the right underlying type. -* `exceptions.hpp`: Provides definitions for all custom exceptions that - might be thrown within FIL and need to be handled by an external - application. -* `postproc_ops.hpp`: Provides enums used to specify how leaf outputs should be - processed. -* `treelite_importer.hpp`: Provides `import_from_treelite_model` and - `import_from_treelite_handle`, either of which can be used to convert a - Treelite model to a `forest_model` object to be used for accelerated - inference. - -### Detail Headers -* `cpu_introspection.hpp`: Provides constants and utilities to evaluate - CPU capabilities for optimized performance. -* `decision_forest_builder.hpp`: Provides generic tools for building - FIL forests from some other source. In the current FIL codebase, the - Treelite import code is the only place this is used, but it could be used - to create import utilities for other sources as well. -* `device_initialization.hpp`: Contains code for anything that must be done - to initialize execution on a device. For GPUs, this may mean setting - specific CUDA options. -* `evaluate_tree.hpp`: Contains code for evaluating a single tree on input - data. -* `forest.hpp`: Provide the storage struct `forest` whose *sole* - responsibility is to hold model data to be used for inference. -* `gpu_introspection.hpp`: Provides constants and utilities to evaluate - GPU capabilities for optimized performance. -* `infer.hpp`: Contains wrapper code for performing inference on a `forest` - object (either on CPU or GPU). This wrapper takes data that has been - extracted from the `forest_model` object if necessary to control details - of forest evaluation. -* `node.hpp`: Provides template for an individual node of a tree. -* `postprocessor.hpp`: Provides device-agnostic code for postprocessing - the output of model leaves. -* `specialization_types.hpp`: Defines all specializations that are used to - construct instantiations of the `decision_forest` template. -* `infer_kernel/`: This directory contains device-specific code that - determines how `evaluate_tree` calls will be performed in parallel. -* `specializations/`: Because there is a large matrix of - specializations for `decision_forest`, it would be tedious and - error-prone to list out all the implementations in source files. - Furthermore, because these templates are complex we wish to avoid - recompiling them unnecessarily. Therefore, this directory contains headers - with macros for declaring the necessary implementations in source files and - declaring the corresponding templates as `extern` elsewhere. Because - these specializations need to be explicitly declared, this must be - implemented as a macro. - -### Source Files -The FIL source files contain no implementation details. They -merely use the macros defined in -`include/cuml/fil/detail/specializations` to indicate the template -instantiations that must be compiled. These are broken up into an arbitrary -number of source files. To improve build parallelization, they could be broken -up further, or to reduce the number of source files, they could be -consolidated. diff --git a/cpp/include/cuml/fil/README.md b/cpp/include/cuml/fil/README.md deleted file mode 100644 index f9da29ec49..0000000000 --- a/cpp/include/cuml/fil/README.md +++ /dev/null @@ -1,161 +0,0 @@ -# Forest Inference Library (FIL) -RAPIDS Forest Inference Library (FIL) provides accelerated inference for -tree-based machine learning models. Unlike packages like XGBoost, -LightGBM, or even Scikit-Learn/cuML's random forest implementations, FIL -cannot be used to _train_ forest models. Instead, its goal is to speed up -inference using forest models trained by all of those packages. - -This directory contains an implementation of FIL which -provides both CPU and GPU execution. Its GPU implementation also offers -improved performance relative to the existing implementation in many but not all cases. - -For Python usage information and more extensive information on -parameter-tuning and other end-user functionality, check out -TODO(wphicks). This document will focus on the C++ implementation, -offering details on both how to use FIL as a library and how to work with it -as a FIL contributor. - -## C++ Usage -All headers required to make use of FIL in another C++ project are -available in the top-level include directory. The `detail` directory -contains implementation details that are not required to use FIL and which -will certainly change over time. - -**A NOTE ON THE `raft_proto` NAMESPACE:** For the first iteration of this FIL -implementation, much of the more general-purpose CPU-GPU interoperable code -has temporarily been put in the `raft_proto` namespace. As the name suggests, -the intention is that most or all of this functionality will either be moved -to RAFT or that RAFT features will be updated to provide CPU-GPU -compatible versions of the same. - -### Importing a model -FIL uses Treelite as a common translation layer for all its input types. -To load a forest model, we first create a Treelite model handle as -follows. Here, we use an XGBoost JSON model as an example, but Treelite has -similar load methods for each of the serialization formats it supports. - -```cpp -auto filename = "xgboost.json"; -auto tl_model = treelite::model_loader::LoadXGBoostModelJSON(filename, "{}"); -``` - -We then import the Treelite model into FIL via the -`import_from_treelite_model` function. All arguments except the first are -optional, but we show them all here for illustration. - -```cpp -auto stream = cudaStream_t{}; -checkCuda(cudaStreamCreate(&stream)); - -auto fil_model = import_from_treelite_model( - *tl_model, // The Treelite model - tree_layout::depth_first, // layout - 128u, // align_bytes - false, // use_double_precision - raft_proto::device_type::gpu, // mem_type - 0, // device_id - stream // CUDA stream -); -``` - -**layout:** The in-memory layout of nodes in the model. Depending on the model, -either `depth_first` or `breadth_first` may offer better performance. -In general, shallow trees benefit from a `breadth_first` layout, and deep trees -benefit from a `depth_first` layout, but this pattern is not absolute. - -**align_bytes:** If given a non-zero value, each tree will be padded to a size -that is a multiple of this value by appending additional empty nodes. This -can offer mild performance benefits by increasing the likelihood that memory -reads begin on a cache line boundary. For GPU execution, a value of 128 is -recommended. For most CPUs, a value of 0 is recommended, although using 64 can -occasionally provide benefits. - -**use_double_precision**: This argument takes a `std::optional`. If -`std::nullopt` is used (the default), the *native* precision of the model -serialization format will be used. Otherwise, the model will be evaluated -at double precision if this value is set to `true` or single precision if this -value is set to `false`. - -**dev_type**: This argument controls where the model will be executed. If `raft_proto::device_type::gpu`, then it will be executed on GPU. If `raft_proto::device_type::cpu`, then it will be executed on CPU. - -**device_id**: This integer indicates the ID of the GPU which should be used. -If CPU is being used, this argument is ignored. - -**stream**: The CUDA stream which will be used for the actual model import. -If CPU is being used, this argument is ignored. Note that you do *not* need -CUDA headers if you are working with a CPU-only build of FIL. This -argument uses a `raft_proto::cuda_stream` type which evaluates to a -placeholder type in CPU-only builds. For applications which themselves want to -implement CPU-GPU interoperable builds, the `raft_proto::cuda_stream` type can be -used directly. - - -### Inference -The `import_from_treelite_model` function will return a `forest_model` object. -This object has several `predict` methods that can be used to return -inference results for the model. We will describe here the one most likely -to be used by external applications: - -```cpp -auto num_rows = std::size_t{1000}; -auto num_outputs = fil_model.num_outputs(); // Outputs per row - -auto output = static_cast(nullptr); // Loaded as single - // precision, so use floats - // for I/O -// Allocate enough space for num_outputs floats per row -cudaMalloc((void**)&output, num_rows * num_outputs * sizeof(float)); - -// Assuming that input is a float* pointing to data already located on-device - -auto handle = raft_proto::handle_t{}; - -fil_model.predict( - handle, - output, - input, - num_rows, - raft_proto::device_type::gpu, // out_mem_type - raft_proto::device_type::gpu, // in_mem_type - 4 // chunk_size -); -``` - -**handle**: To provide a unified interface on CPU and GPU, we introduce -`raft_proto::handle_t` as a wrapper for `raft::handle_t`. This is currently just a -placeholder in CPU-only builds, and using it does not require any CUDA -functionality. For GPU-enabled builds, you can construct a -`raft_proto_handle_t` directly from the `raft::handle_t` you wish to use. - -**output**: Pointer to pre-allocated buffer where results should be -written. If the model has been loaded at single precision, this should be a -`float*`. If the model has been loaded at double precision, this should be a -`double*`. - -**input**: Pointer to the input data (in C-major order). If the model has been -loaded at single precision, this should be a `float*`. If the model has been -loaded at double precision, this should be a `double*`. - -**num_rows**: The number of input rows. - -**out_mem_type**: Indicates whether output buffer is on device or host. - -**in_mem_type**: Indicates whether input buffer is on device or host. - -**chunk_size**: This value has a somewhat different meaning for CPU and GPU, -but it generally indicates the number of rows which are evaluated in a single -iteration of FIL's forest evaluation algorithm. On GPU, any power of 2 from 1 to 32 -may be used for this value, and *in general* larger batches benefit from -higher values. Optimizing this value can make an *enormous* difference -in performance and depends on both the model and hardware used to run it. On -CPU, this parameter can take on any value, but powers of 2 between 1 and 512 -are recommended. A default value of 64 is generally a safe choice, unless the -batch size is less than 64, in which case a smaller value is recommended. In -general, larger batch sizes benefit from higher chunk size values. This -argument is a `std::optional`, and if `std::nullopt` is passed, a chunk size -will be selected based on heuristics. - -## Learning More -While the above usage summary should be enough to get started using FIL in -another C++ application, you can learn more about the details of this -implementation by reading TODO(wphicks). diff --git a/cpp/include/cuml/fil/constants.hpp b/cpp/include/cuml/fil/constants.hpp deleted file mode 100644 index e074607448..0000000000 --- a/cpp/include/cuml/fil/constants.hpp +++ /dev/null @@ -1,27 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -namespace ML { -namespace fil { -/** - * The default memory layout for FIL trees if not otherwise specified - */ -auto constexpr static const preferred_tree_layout = tree_layout::breadth_first; -/** - * The number of bits used for flags in node metadata - * - * Each node in a FIL tree must specify the feature used for its split in - * addition to some other basic information. The feature ID is "packed" - * with a few flags in order to reduce the size of the node. This constant - * indicates how many leading bits are reserved for flags to allow import - * functions to assess how much space is required for the whole metadata - * field. - */ -auto constexpr static const reserved_node_metadata_bits = 3; - -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/decision_forest.hpp b/cpp/include/cuml/fil/decision_forest.hpp deleted file mode 100644 index a18322aefe..0000000000 --- a/cpp/include/cuml/fil/decision_forest.hpp +++ /dev/null @@ -1,483 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include - -#include -#include -#include -#include -#include - -namespace ML { -namespace fil { - -/** - * A general-purpose decision forest implementation - * - * This template provides an optimized but generic implementation of a decision - * forest. Template parameters are used to specialize the - * implementation based on the size and characteristics of the forest. - * For instance, the smallest integer that can express the offset between a - * parent and child node within a tree is used in order to minimize the size - * of a node, increasing the number that can fit within the L2 or L1 cache. - * - * @tparam layout_v The in-memory layout of nodes in this forest - * @tparam threshold_t The floating-point type used for quantitative splits - * @tparam index_t The integer type used for storing many things within a - * forest, including the category value of categorical nodes and the index at - * which vector output for a leaf node is stored. - * @tparam metadata_storage_t The type used for storing node metadata. - * The first several bits will be used to store flags indicating various - * characteristics of the node, and the remaining bits provide the integer - * index of the feature for this node's split - * @tparam offset_t An integer used to indicate the offset between a node and - * its most distant child. This type must be large enough to store the - * largest such offset in the entire forest. - */ -template -struct decision_forest { - /** - * The in-memory layout of nodes in this forest - */ - auto constexpr static const layout = layout_v; - /** - * The type of the forest object which is actually passed to the CPU/GPU - * for inference - */ - using forest_type = forest; - /** - * The type of nodes within the forest - */ - using node_type = typename forest_type::node_type; - /** - * The type used for input and output to the model - */ - using io_type = typename forest_type::io_type; - /** - * The type used for quantitative splits within the model - */ - using threshold_type = threshold_t; - /** - * The type used to indicate how leaf output should be post-processed - */ - using postprocessor_type = postprocessor; - /** - * The type used for storing data on categorical nodes - */ - using categorical_storage_type = typename node_type::index_type; - - /** - * Construct an empty decision forest - */ - decision_forest() - : nodes_{}, - root_node_indexes_{}, - node_id_mapping_{}, - bias_{}, - vector_output_{}, - categorical_storage_{}, - num_features_{}, - num_outputs_{}, - leaf_size_{}, - has_categorical_nodes_{false}, - row_postproc_{}, - elem_postproc_{}, - average_factor_{}, - postproc_constant_{} - { - } - - /** - * Construct a decision forest with the indicated data - * - * @param nodes A buffer containing all nodes within the forest - * @param root_node_indexes A buffer containing the index of the root node - * of every tree in the forest - * @param node_id_mapping Mapping to use to convert FIL's internal node ID into Treelite's node - * ID. Only relevant when predict_type == infer_kind::leaf_id - * @param bias The bias term that is added to the output as part of the postprocessing step. - * The bias term should have same length as num_outputs. - * @param num_features The number of features per input sample for this model - * @param num_outputs The number of outputs per row from this model - * @param has_categorical_nodes Whether this forest contains any - * categorical nodes - * @param vector_output A buffer containing the output from all vector - * leaves for this model. Each leaf node will specify the offset within - * this buffer at which its vector output begins, and leaf_size will be - * used to determine how many subsequent entries from the buffer should be - * used to construct the vector output. A value of std::nullopt indicates - * that this is not a vector leaf model. - * @param categorical_storage For models with inputs on too many categories - * to be stored in the bits of an `index_t`, it may be necessary to store - * categorical information external to the node itself. This buffer - * contains the necessary storage for this information. - * @param leaf_size The number of output values per leaf (1 for non-vector - * leaves; >1 for vector leaves) - * @param row_postproc The post-processing operation to be applied to an - * entire row of the model output - * @param elem_postproc The per-element post-processing operation to be - * applied to the model output - * @param average_factor A factor which is used for output - * normalization - * @param postproc_constant A constant used by some post-processing - * operations, including sigmoid, exponential, and - * logarithm_one_plus_exp - */ - decision_forest(raft_proto::buffer&& nodes, - raft_proto::buffer&& root_node_indexes, - raft_proto::buffer&& node_id_mapping, - raft_proto::buffer&& bias, - index_type num_features, - index_type num_outputs = index_type{2}, - bool has_categorical_nodes = false, - std::optional>&& vector_output = std::nullopt, - std::optional>&& - categorical_storage = std::nullopt, - index_type leaf_size = index_type{1}, - row_op row_postproc = row_op::disable, - element_op elem_postproc = element_op::disable, - io_type average_factor = io_type{1}, - io_type postproc_constant = io_type{1}) - : nodes_{nodes}, - root_node_indexes_{root_node_indexes}, - node_id_mapping_{node_id_mapping}, - bias_{bias}, - vector_output_{vector_output}, - categorical_storage_{categorical_storage}, - num_features_{num_features}, - num_outputs_{num_outputs}, - leaf_size_{leaf_size}, - has_categorical_nodes_{has_categorical_nodes}, - row_postproc_{row_postproc}, - elem_postproc_{elem_postproc}, - average_factor_{average_factor}, - postproc_constant_{postproc_constant} - { - if (nodes.memory_type() != root_node_indexes.memory_type()) { - throw raft_proto::mem_type_mismatch( - "Nodes and indexes of forest must both be stored on either host or device"); - } - if (nodes.device_index() != root_node_indexes.device_index()) { - throw raft_proto::mem_type_mismatch( - "Nodes and indexes of forest must both be stored on same device"); - } - detail::initialize_device(nodes.device()); - } - - /** The number of features per row expected by the model */ - auto num_features() const { return num_features_; } - /** The number of trees in the model */ - auto num_trees() const { return root_node_indexes_.size(); } - /** Whether or not leaf nodes have vector outputs */ - auto has_vector_leaves() const { return vector_output_.has_value(); } - - /** - * The number of outputs per row generated by the model for the given - * type of inference. Note: This will differ from num_outputs argument - * passed to the constructor, if inference_kind is not default_kind. - */ - auto num_outputs(infer_kind inference_kind = infer_kind::default_kind) const - { - auto result = num_outputs_; - if (inference_kind == infer_kind::per_tree) { - result = num_trees(); - if (has_vector_leaves()) { result *= num_outputs_; } - } else if (inference_kind == infer_kind::leaf_id) { - result = num_trees(); - } - return result; - } - - /** The operation used for postprocessing all outputs for a single row */ - auto row_postprocessing() const { return row_postproc_; } - // Setter for row_postprocessing - void set_row_postprocessing(row_op val) { row_postproc_ = val; } - /** The operation used for postprocessing each element of the output for a - * single row */ - auto elem_postprocessing() const { return elem_postproc_; } - - /** The type of memory (device/host) where the model is stored */ - auto memory_type() { return nodes_.memory_type(); } - /** The ID of the device on which this model is loaded */ - auto device_index() { return nodes_.device_index(); } - - /** - * Perform inference with this model - * - * @param[out] output The buffer where the model output should be stored. - * This must be of size ROWS x num_outputs(). - * @param[in] input The buffer containing the input data - * @param[in] stream For GPU execution, the CUDA stream. For CPU execution, - * this optional parameter can be safely omitted. - * @param[in] predict_type Type of inference to perform. Defaults to summing - * the outputs of all trees and produce an output per row. If set to - * "per_tree", we will instead output all outputs of individual trees. - * If set to "leaf_id", we will output the integer ID of the leaf node - * for each tree. - * @param[in] specified_rows_per_block_iter If non-nullopt, this value is - * used to determine how many rows are evaluated for each inference - * iteration within a CUDA block. Runtime performance is quite sensitive - * to this value, but it is difficult to predict a priori, so it is - * recommended to perform a search over possible values with realistic - * batch sizes in order to determine the optimal value. Any power of 2 from - * 1 to 32 is a valid value, and in general larger batches benefit from - * larger values. - */ - void predict(raft_proto::buffer& output, - raft_proto::buffer const& input, - raft_proto::cuda_stream stream = raft_proto::cuda_stream{}, - infer_kind predict_type = infer_kind::default_kind, - std::optional specified_rows_per_block_iter = std::nullopt) - { - if (output.memory_type() != memory_type() || input.memory_type() != memory_type()) { - throw raft_proto::wrong_device_type{ - "Tried to use host I/O data with model on device or vice versa"}; - } - if (output.device_index() != device_index() || input.device_index() != device_index()) { - throw raft_proto::wrong_device{"I/O data on different device than model"}; - } - auto* vector_output_data = - (vector_output_.has_value() ? vector_output_->data() : static_cast(nullptr)); - auto* categorical_storage_data = - (categorical_storage_.has_value() ? categorical_storage_->data() - : static_cast(nullptr)); - switch (nodes_.device().index()) { - case 0: - fil::detail::infer(obj(), - get_postprocessor(predict_type), - output.data(), - input.data(), - index_type(input.size() / num_features_), - num_features_, - num_outputs(predict_type), - has_categorical_nodes_, - vector_output_data, - categorical_storage_data, - predict_type, - specified_rows_per_block_iter, - std::get<0>(nodes_.device()), - stream); - break; - case 1: - fil::detail::infer(obj(), - get_postprocessor(predict_type), - output.data(), - input.data(), - index_type(input.size() / num_features_), - num_features_, - num_outputs(predict_type), - has_categorical_nodes_, - vector_output_data, - categorical_storage_data, - predict_type, - specified_rows_per_block_iter, - std::get<1>(nodes_.device()), - stream); - break; - } - } - - private: - /** The nodes for all trees in the forest */ - raft_proto::buffer nodes_; - /** The index of the root node for each tree in the forest */ - raft_proto::buffer root_node_indexes_; - /** Mapping to apply to node IDs. Only relevant when predict_type == infer_kind::leaf_id */ - raft_proto::buffer node_id_mapping_; - /** Bias term to apply to the output */ - raft_proto::buffer bias_; - /** Buffer of outputs for all leaves in vector-leaf models */ - std::optional> vector_output_; - /** Buffer of elements used as backing data for bitsets which specify - * categories for all categorical nodes in the model. */ - std::optional> categorical_storage_; - - // Metadata - index_type num_features_; - index_type num_outputs_; - index_type leaf_size_; - bool has_categorical_nodes_ = false; - // Postprocessing constants - row_op row_postproc_; - element_op elem_postproc_; - io_type average_factor_; - io_type postproc_constant_; - - auto obj() const - { - return forest_type{nodes_.data(), - root_node_indexes_.data(), - node_id_mapping_.data(), - bias_.data(), - static_cast(root_node_indexes_.size()), - num_outputs_}; - } - - auto get_postprocessor(infer_kind inference_kind = infer_kind::default_kind) const - { - auto result = postprocessor_type{}; - if (inference_kind == infer_kind::default_kind) { - result = - postprocessor_type{row_postproc_, elem_postproc_, average_factor_, postproc_constant_}; - } - return result; - } - - auto leaf_size() const { return leaf_size_; } -}; - -namespace detail { -/** - * A convenience wrapper to simplify template instantiation of - * decision_forest - * - * This template takes the large range of available template parameters - * and reduces them to just three standard choices. - * - * @tparam layout The in-memory layout of nodes in this forest - * @tparam double_precision Whether this model should use double-precision - * for floating-point evaluation and 64-bit integers for indexes - * @tparam large_trees Whether this forest expects more than 2**(16 -3) - 1 = - * 8191 features or contains nodes whose child is offset more than 2**16 - 1 = 65535 nodes away. - */ -template -using preset_decision_forest = decision_forest< - layout, - typename specialization_types::threshold_type, - typename specialization_types::index_type, - typename specialization_types::metadata_type, - typename specialization_types::offset_type>; - -} // namespace detail - -/** A variant containing all standard decision_forest instantiations */ -using decision_forest_variant = std::variant< - detail::preset_decision_forest< - std::variant_alternative_t<0, detail::specialization_variant>::layout, - std::variant_alternative_t<0, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<0, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<1, detail::specialization_variant>::layout, - std::variant_alternative_t<1, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<1, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<2, detail::specialization_variant>::layout, - std::variant_alternative_t<2, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<2, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<3, detail::specialization_variant>::layout, - std::variant_alternative_t<3, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<3, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<4, detail::specialization_variant>::layout, - std::variant_alternative_t<4, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<4, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<5, detail::specialization_variant>::layout, - std::variant_alternative_t<5, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<5, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<6, detail::specialization_variant>::layout, - std::variant_alternative_t<6, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<6, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<7, detail::specialization_variant>::layout, - std::variant_alternative_t<7, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<7, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<8, detail::specialization_variant>::layout, - std::variant_alternative_t<8, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<8, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<9, detail::specialization_variant>::layout, - std::variant_alternative_t<9, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<9, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<10, detail::specialization_variant>::layout, - std::variant_alternative_t<10, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<10, detail::specialization_variant>::has_large_trees>, - detail::preset_decision_forest< - std::variant_alternative_t<11, detail::specialization_variant>::layout, - std::variant_alternative_t<11, detail::specialization_variant>::is_double_precision, - std::variant_alternative_t<11, detail::specialization_variant>::has_large_trees>>; - -/** - * Determine the variant index of the decision_forest type to used based on - * model characteristics - * - * @param use_double_thresholds Whether single or double-precision floating - * point values should be used for quantitative splits - * @param max_node_offset The largest offset between a parent node and either - * of its children - * @param num_features The number of input features per row - * @param num_categorical_nodes The total number of categorical nodes in the - * forest - * @param max_num_categories The maximum number of categories in any - * categorical feature used by the model - * @param num_vector_leaves The total number of leaf nodes which produce vector - * outputs. For non-vector-leaf models, this should be 0. For vector-leaf - * models, this should be the total number of leaf nodes. - * @param layout The in-memory layout to be used for nodes in the forest - */ -inline auto get_forest_variant_index(bool use_double_thresholds, - index_type max_node_offset, - index_type num_features, - index_type num_categorical_nodes = index_type{}, - index_type max_num_categories = index_type{}, - index_type num_vector_leaves = index_type{}, - tree_layout layout = preferred_tree_layout) -{ - using small_index_t = - typename detail::specialization_types::index_type; - auto max_local_categories = index_type(sizeof(small_index_t) * 8); - // If the index required for pointing to categorical storage bins or vector - // leaf output exceeds what we can store in a uint32_t, uint64_t will be used - // - // TODO(wphicks): We are overestimating categorical storage required here - auto double_indexes_required = - (max_num_categories > max_local_categories && - ((raft_proto::ceildiv(max_num_categories, max_local_categories) + 1 * num_categorical_nodes) > - std::numeric_limits::max())) || - num_vector_leaves > std::numeric_limits::max(); - - auto double_precision = use_double_thresholds || double_indexes_required; - - using small_metadata_t = - typename detail::specialization_types::metadata_type; - using small_offset_t = - typename detail::specialization_types::offset_type; - - auto large_trees = - (num_features > (std::numeric_limits::max() >> reserved_node_metadata_bits) || - max_node_offset > std::numeric_limits::max()); - - auto layout_value = static_cast>(layout); - - return ((index_type{layout_value} << index_type{2}) + - (index_type{double_precision} << index_type{1}) + index_type{large_trees}); -} -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/bitset.hpp b/cpp/include/cuml/fil/detail/bitset.hpp deleted file mode 100644 index edbe5140ed..0000000000 --- a/cpp/include/cuml/fil/detail/bitset.hpp +++ /dev/null @@ -1,110 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include - -#include -#include -#include - -#ifndef __CUDACC__ -#include -#endif - -namespace ML { -namespace fil { -namespace detail { -template -struct bitset { - using storage_type = storage_t; - using index_type = index_t; - - // Ensrue that index_t is unsigned. Bound checks below rely on index_t being unsigned - static_assert(std::is_unsigned_v, "index_t must be unsigned"); - - auto constexpr static const bin_width = index_type(sizeof(storage_type) * 8); - - HOST DEVICE bitset() : data_{nullptr}, num_bits_{0} {} - - HOST DEVICE bitset(storage_type* data, index_type size) : data_{data}, num_bits_{size} {} - - HOST DEVICE bitset(storage_type* data) : data_{data}, num_bits_(sizeof(storage_type) * 8) {} - - HOST DEVICE auto size() const { return num_bits_; } - HOST DEVICE auto bin_count() const - { - return num_bits_ / bin_width + (num_bits_ % bin_width != 0); - } - - // Standard bit-wise mutators and accessor - HOST DEVICE auto& set(index_type index) - { - // Guard against OOB writes; silently ignored to preserve memory safety - if (index < num_bits_) { data_[bin_from_index(index)] |= mask_in_bin(index); } - return *this; - } - HOST DEVICE auto& clear(index_type index) - { - if (index < num_bits_) { data_[bin_from_index(index)] &= ~mask_in_bin(index); } - return *this; - } - HOST DEVICE auto test(index_type index) const - { - auto result = false; - if (index < num_bits_) { result = ((data_[bin_from_index(index)] & mask_in_bin(index)) != 0); } - return result; - } - HOST DEVICE auto& flip() - { - for (auto i = index_type{}; i < bin_count(); ++i) { - data_[i] = ~data_[i]; - } - return *this; - } - - // Bit-wise boolean operations - HOST DEVICE auto& operator&=(bitset const& other) - { - for (auto i = index_type{}; i < min(size(), other.size()); ++i) { - data_[i] &= other.data_[i]; - } - return *this; - } - HOST DEVICE auto& operator|=(bitset const& other) - { - for (auto i = index_type{}; i < min(size(), other.size()); ++i) { - data_[i] |= other.data_[i]; - } - return *this; - } - HOST DEVICE auto& operator^=(bitset const& other) - { - for (auto i = index_type{}; i < min(size(), other.size()); ++i) { - data_[i] ^= other.data_[i]; - } - return *this; - } - HOST DEVICE auto& operator~() const - { - flip(); - return *this; - } - - private: - storage_type* data_; - index_type num_bits_; - - HOST DEVICE auto mask_in_bin(index_type index) const - { - return storage_type{1} << (index % bin_width); - } - - HOST DEVICE auto bin_from_index(index_type index) const { return index / bin_width; } -}; - -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/cpu_introspection.hpp b/cpp/include/cuml/fil/detail/cpu_introspection.hpp deleted file mode 100644 index b1e5936c28..0000000000 --- a/cpp/include/cuml/fil/detail/cpu_introspection.hpp +++ /dev/null @@ -1,19 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include - -namespace ML { -namespace fil { -namespace detail { -#ifdef __cpplib_hardware_interference_size -using std::hardware_constructive_interference_size; -#else -auto constexpr static const hardware_constructive_interference_size = std::size_t{64}; -#endif -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/decision_forest_builder.hpp b/cpp/include/cuml/fil/detail/decision_forest_builder.hpp deleted file mode 100644 index 488e1be893..0000000000 --- a/cpp/include/cuml/fil/detail/decision_forest_builder.hpp +++ /dev/null @@ -1,361 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -namespace ML { -namespace fil { -namespace detail { - -struct floating_point_truncation_error : std::exception { - floating_point_truncation_error() {} - floating_point_truncation_error(std::string msg) : msg_{msg} {} - floating_point_truncation_error(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_.c_str(); } - - private: - std::string msg_; -}; - -template -To safe_cast_floating_point(From x) -{ - static_assert(std::is_floating_point_v && std::is_floating_point_v, - "Source and destination types must be both floating-point types."); - if constexpr (sizeof(To) >= sizeof(From)) { - // Widening cast - return static_cast(x); - } else { - // Narrowing cast: should be checked - if (!std::isfinite(x)) { - throw floating_point_truncation_error{"Cannot cast an INF or NaN value"}; - } - auto constexpr lower_limit = static_cast(std::numeric_limits::lowest()); - auto constexpr upper_limit = static_cast(std::numeric_limits::max()); - if (x < lower_limit) { - std::ostringstream ss; - ss << "Input must be at least " << lower_limit << "."; - throw floating_point_truncation_error{ss.str()}; - } - if (x > upper_limit) { - std::ostringstream ss; - ss << "Input must be at most " << upper_limit << "."; - throw floating_point_truncation_error{ss.str()}; - } - return static_cast(x); - } -} - -/* - * Struct used to build FIL forests - */ -template -struct decision_forest_builder { - /* The type for nodes in the given decision_forest type */ - using node_type = typename decision_forest_t::node_type; - - /* Add a node with a categorical split */ - template - void add_categorical_node( - iter_t vec_begin, - iter_t vec_end, - std::optional tl_node_id = std::nullopt, - std::size_t depth = std::size_t{1}, - bool default_to_distant_child = false, - typename node_type::metadata_storage_type feature = typename node_type::metadata_storage_type{}, - typename node_type::offset_type offset = typename node_type::offset_type{}) - { - auto constexpr const bin_width = - typename node_type::index_type{sizeof(typename node_type::index_type) * 8}; - auto node_value = typename node_type::index_type{}; - auto set_storage = &node_value; - - // Check invariants for data types - using cat_t = typename std::iterator_traits::value_type; - using index_t = typename node_type::index_type; - static_assert(std::is_same_v, "Category value must be uint32_t"); - static_assert(std::is_same_v || std::is_same_v, - "Index type in tree node must be either uint32_t or uint64_t"); - - // Ensure that (max_cat + 1) can be represented as index_t to prevent integer overflow. - auto max_cat = (vec_begin != vec_end) ? *std::max_element(vec_begin, vec_end) : cat_t{0}; - if constexpr (std::is_same_v) { - if (max_cat == std::numeric_limits::max()) { - throw model_import_error{std::string{"Category index must be at most "} + - std::to_string(std::numeric_limits::max() - 1)}; - } - } - auto max_cat_plus_one = static_cast(max_cat) + index_t{1}; - - if (max_num_categories_ > bin_width) { - node_value = categorical_storage_.size(); - auto bins_required = raft_proto::ceildiv(max_cat_plus_one, bin_width); - categorical_storage_.push_back(max_cat_plus_one); - categorical_storage_.resize(categorical_storage_.size() + bins_required); - set_storage = &(categorical_storage_[node_value + 1]); - } - auto set = bitset{set_storage, max_cat_plus_one}; - std::for_each(vec_begin, vec_end, [&set](auto&& cat_index) { set.set(cat_index); }); - - add_node( - node_value, tl_node_id, depth, false, default_to_distant_child, true, feature, offset, false); - } - - /* Add a leaf node with vector output */ - template - void add_leaf_vector_node(iter_t vec_begin, - iter_t vec_end, - std::optional tl_node_id = std::nullopt, - std::size_t depth = std::size_t{1}) - { - auto leaf_index = typename node_type::index_type(vector_output_.size() / output_size_); - std::copy(vec_begin, vec_end, std::back_inserter(vector_output_)); - - add_node(leaf_index, - tl_node_id, - depth, - true, - false, - false, - typename node_type::metadata_storage_type{}, - typename node_type::offset_type{}, - false); - } - - /* Add a node to the model */ - template - void add_node( - value_t val, - std::optional tl_node_id = std::nullopt, - std::size_t depth = std::size_t{1}, - bool is_leaf_node = true, - bool default_to_distant_child = false, - bool is_categorical_node = false, - typename node_type::metadata_storage_type feature = typename node_type::metadata_storage_type{}, - typename node_type::offset_type offset = typename node_type::offset_type{}, - bool is_inclusive = false) - { - if (depth == std::size_t{}) { - if (alignment_ != index_type{}) { - if (cur_node_index_ % alignment_ != index_type{}) { - auto padding = (alignment_ - cur_node_index_ % alignment_); - for (auto i = index_type{}; i < padding; ++i) { - add_node(typename node_type::threshold_type{}, std::nullopt); - } - } - } - root_node_indexes_.push_back(cur_node_index_); - } - - if (is_inclusive) { val = std::nextafter(val, std::numeric_limits::infinity()); } - nodes_.emplace_back( - val, is_leaf_node, default_to_distant_child, is_categorical_node, feature, offset); - // 0 indicates the lack of ID mapping for a particular node - node_id_mapping_.push_back(static_cast(tl_node_id.value_or(0))); - ++cur_node_index_; - } - - /* Set the element-wise postprocessing operation for this model */ - void set_element_postproc(element_op val) { element_postproc_ = val; } - /* Set the row-wise postprocessing operation for this model */ - void set_row_postproc(row_op val) { row_postproc_ = val; } - /* Set the value to divide by during postprocessing */ - void set_average_factor(double val) { average_factor_ = val; } - /* Set the bias term, which is added to the output. The bias term - * should have the same length as output_size. */ - void set_bias(std::vector val) - { - bias_.resize(val.size()); - std::transform(val.begin(), val.end(), bias_.begin(), [](double e) { - return static_cast(e); - }); - } - /* Set the value of the constant used in the postprocessing operation - * (if any) */ - void set_postproc_constant(double val) { postproc_constant_ = val; } - /* Set the number of outputs per row for this model */ - void set_output_size(index_type val) - { - if (output_size_ != index_type{1} && output_size_ != val) { - throw unusable_model_exception("Inconsistent leaf vector size"); - } - output_size_ = val; - } - - decision_forest_builder(index_type max_num_categories = index_type{}, - index_type align_bytes = index_type{}) - : cur_node_index_{}, - max_num_categories_{max_num_categories}, - alignment_{std::lcm(align_bytes, index_type(sizeof(node_type)))}, - output_size_{1}, - row_postproc_{}, - element_postproc_{}, - average_factor_{}, - postproc_constant_{}, - nodes_{}, - root_node_indexes_{}, - vector_output_{}, - bias_{} - { - } - - /* Return the FIL decision forest built by this builder */ - auto get_decision_forest(index_type num_feature, - index_type num_class, - raft_proto::device_type mem_type = raft_proto::device_type::cpu, - int device = 0, - raft_proto::cuda_stream stream = raft_proto::cuda_stream{}) - { - // Validate forest invariants the inference kernel relies on. After this - // function returns, the forest is treated as trusted by the kernel. - - // tree_index arithmetic in the kernel uses index_type, so the tree count - // must fit without narrowing. - if (root_node_indexes_.size() > std::numeric_limits::max()) { - throw model_import_error{std::string{"Forest has "} + - std::to_string(root_node_indexes_.size()) + - " trees, which exceeds the maximum representable in index_type (" + - std::to_string(std::numeric_limits::max()) + ")"}; - } - - // forest::get_tree_root(tree_index) dereferences nodes_ + root_index. - // Ensure each root index points into the nodes buffer. - for (auto i = std::size_t{0}; i < root_node_indexes_.size(); ++i) { - if (root_node_indexes_[i] >= nodes_.size()) { - throw model_import_error{ - std::string{"Tree "} + std::to_string(i) + ": root node index out of bounds (" + - std::to_string(root_node_indexes_[i]) + " >= " + std::to_string(nodes_.size()) + ")"}; - } - } - - auto constexpr const cat_bin_width = - typename node_type::index_type{sizeof(typename node_type::index_type) * 8}; - if (max_num_categories_ > cat_bin_width) { - auto const storage_size = categorical_storage_.size(); - for (auto i = std::size_t{0}; i < nodes_.size(); ++i) { - auto const& n = nodes_[i]; - if (n.is_leaf() || !n.is_categorical()) { continue; } - auto const offset = n.index(); - - // evaluate_tree_impl() reads categorical_storage[offset] as the number - // of categories for this node; offset must be in-range. - if (offset >= storage_size) { - throw model_import_error{std::string{"Categorical node "} + std::to_string(i) + - ": storage offset out of bounds (" + std::to_string(offset) + - " >= " + std::to_string(storage_size) + ")"}; - } - auto const stored_num_cats = categorical_storage_[offset]; - auto const bins_required = raft_proto::ceildiv(stored_num_cats, cat_bin_width); - - // evaluate_tree_impl() reconstructs a bitset from - // [offset + 1, offset + 1 + bins_required). Compute this range using - // size_t to keep the arithmetic explicit and overflow-safe. - auto const bits_begin = static_cast(offset) + std::size_t{1}; - auto const bits_end = bits_begin + static_cast(bins_required); - if (bits_end > storage_size) { - throw model_import_error{std::string{"Categorical node "} + std::to_string(i) + - ": bitset extends past categorical_storage end"}; - } - } - } - - // Safely cast average_factor_ and postproc_constant_ to node_type::threshold_type - auto average_factor_casted = typename node_type::threshold_type{}; - auto postproc_constant_casted = typename node_type::threshold_type{}; - try { - average_factor_casted = - safe_cast_floating_point(average_factor_); - // We can't use cuda::narrow here, because it throws for imprecise conversion, i.e. casting - // double{3.1} to float. - } catch (const floating_point_truncation_error& e) { - throw unusable_model_exception{std::string{"Found an invalid value for averaging factor: "} + - e.what()}; - } - try { - postproc_constant_casted = - safe_cast_floating_point(postproc_constant_); - } catch (const floating_point_truncation_error& e) { - throw unusable_model_exception{ - std::string{"Found an invalid value for postprocessing constant: "} + e.what()}; - } - return decision_forest_t{ - raft_proto::buffer{ - raft_proto::buffer{nodes_.data(), nodes_.size()}, mem_type, device, stream}, - raft_proto::buffer{raft_proto::buffer{root_node_indexes_.data(), root_node_indexes_.size()}, - mem_type, - device, - stream}, - raft_proto::buffer{raft_proto::buffer{node_id_mapping_.data(), node_id_mapping_.size()}, - mem_type, - device, - stream}, - raft_proto::buffer{raft_proto::buffer{bias_.data(), bias_.size()}, mem_type, device, stream}, - num_feature, - num_class, - max_num_categories_ != 0, - vector_output_.empty() - ? std::nullopt - : std::make_optional>( - raft_proto::buffer{vector_output_.data(), vector_output_.size()}, - mem_type, - device, - stream), - categorical_storage_.empty() - ? std::nullopt - : std::make_optional>( - raft_proto::buffer{categorical_storage_.data(), categorical_storage_.size()}, - mem_type, - device, - stream), - output_size_, - row_postproc_, - element_postproc_, - average_factor_casted, - postproc_constant_casted}; - } - - private: - index_type cur_node_index_; - index_type max_num_categories_; - index_type alignment_; - index_type output_size_; - row_op row_postproc_; - element_op element_postproc_; - double average_factor_; - double postproc_constant_; - - std::vector nodes_; - std::vector root_node_indexes_; - std::vector vector_output_; - std::vector bias_; - std::vector categorical_storage_; - std::vector node_id_mapping_; -}; - -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/degenerate_trees.hpp b/cpp/include/cuml/fil/detail/degenerate_trees.hpp deleted file mode 100644 index 5dcf19d208..0000000000 --- a/cpp/include/cuml/fil/detail/degenerate_trees.hpp +++ /dev/null @@ -1,75 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include - -#include - -#include -#include -#include - -namespace ML::fil::detail { - -// This function returns a modified copy of a given Treelite model if it contains -// at least one degenerate tree (a single root node with no child). -// If the model contains no degenerate tree, then the function returns nullptr. -inline std::unique_ptr convert_degenerate_trees(treelite::Model const& tl_model) -{ - bool contains_degenerate = - ML::forest::tree_accumulate(tl_model, false, [](auto&& contains, auto&& tree) { - return contains || tree.IsLeaf(ML::forest::TREELITE_NODE_ID_T{}); - }); - - if (contains_degenerate) { - // Make a copy of the Treelite model, and then update the trees in-place - auto modified_model = treelite::ConcatenateModelObjects({&tl_model}); - std::visit( - [](auto&& concrete_tl_model) { - using model_t = std::remove_const_t>; - using tree_t = - treelite::Tree; - auto modified_trees = std::vector{}; - const auto root_id = ML::forest::TREELITE_NODE_ID_T{}; - for (tree_t& tree : concrete_tl_model.trees) { - if (tree.IsLeaf(root_id)) { - const auto inst_cnt = - tree.HasDataCount(root_id) ? tree.DataCount(root_id) : std::uint64_t{}; - auto new_tree = tree_t{}; - new_tree.Init(); - const auto root_id = new_tree.AllocNode(); - const auto cleft_id = new_tree.AllocNode(); - const auto cright_id = new_tree.AllocNode(); - new_tree.SetChildren(root_id, cleft_id, cright_id); - new_tree.SetNumericalTest( - root_id, int{}, typename model_t::threshold_type{}, true, treelite::Operator::kLE); - if (tree.HasLeafVector(root_id)) { - const auto leaf_vector = tree.LeafVector(root_id); - new_tree.SetLeafVector(cleft_id, leaf_vector); - new_tree.SetLeafVector(cright_id, leaf_vector); - } else { - const auto leaf_value = tree.LeafValue(root_id); - new_tree.SetLeaf(cleft_id, leaf_value); - new_tree.SetLeaf(cright_id, leaf_value); - } - new_tree.SetDataCount(root_id, inst_cnt); - new_tree.SetDataCount(cleft_id, inst_cnt); - new_tree.SetDataCount(cright_id, std::uint64_t{}); - modified_trees.push_back(std::move(new_tree)); - } else { - modified_trees.push_back(std::move(tree)); - } - } - concrete_tl_model.trees = std::move(modified_trees); - }, - modified_model->variant_); - return modified_model; - } else { - return std::unique_ptr(); - } -} - -} // namespace ML::fil::detail diff --git a/cpp/include/cuml/fil/detail/device_initialization.hpp b/cpp/include/cuml/fil/detail/device_initialization.hpp deleted file mode 100644 index 0c34c8dc68..0000000000 --- a/cpp/include/cuml/fil/detail/device_initialization.hpp +++ /dev/null @@ -1,36 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include - -#include -#ifdef CUML_ENABLE_GPU -#include -#endif - -namespace ML { -namespace fil { -namespace detail { -/* Set any required device options for optimizing FIL compute */ -template -void initialize_device(raft_proto::device_id device) -{ - device_initialization::initialize_device(device); -} - -/* Set any required device options for optimizing FIL compute */ -template -void initialize_device(raft_proto::device_id_variant device) -{ - std::visit( - [](auto&& concrete_device) { - device_initialization::initialize_device(concrete_device); - }, - device); -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/device_initialization/cpu.hpp b/cpp/include/cuml/fil/detail/device_initialization/cpu.hpp deleted file mode 100644 index 4a40103f70..0000000000 --- a/cpp/include/cuml/fil/detail/device_initialization/cpu.hpp +++ /dev/null @@ -1,34 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include -#include -#include - -#include - -namespace ML { -namespace fil { -namespace detail { -namespace device_initialization { - -/* Specialization for any initialization required for CPUs - * - * This specialization will also be used for non-GPU-enabled builds - * (as a GPU no-op). - */ -template -std::enable_if_t, - std::bool_constant>, - void> -initialize_device(raft_proto::device_id device) -{ -} - -} // namespace device_initialization -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/device_initialization/gpu.cuh b/cpp/include/cuml/fil/detail/device_initialization/gpu.cuh deleted file mode 100644 index 0287186fac..0000000000 --- a/cpp/include/cuml/fil/detail/device_initialization/gpu.cuh +++ /dev/null @@ -1,242 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include - -namespace ML { -namespace fil { -namespace detail { -namespace device_initialization { - -/* The implementation of the template used to initialize GPU device options - * - * On GPU-enabled builds, the GPU specialization of this template ensures that - * the inference kernels have access to the maximum available dynamic shared - * memory. - */ -template -std::enable_if_t, - std::bool_constant>, - void> -initialize_device(raft_proto::device_id device) -{ - auto device_context = raft_proto::device_setter(device); - auto max_shared_mem_per_block = get_max_shared_mem_per_block(device); - // Run solely for side-effect of caching SM count - get_sm_count(device); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check(cudaFuncSetAttribute( - infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); - raft_proto::cuda_check( - cudaFuncSetAttribute(infer_kernel, - cudaFuncAttributeMaxDynamicSharedMemorySize, - max_shared_mem_per_block)); -} - -CUML_FIL_INITIALIZE_DEVICE(extern template, 0) -CUML_FIL_INITIALIZE_DEVICE(extern template, 1) -CUML_FIL_INITIALIZE_DEVICE(extern template, 2) -CUML_FIL_INITIALIZE_DEVICE(extern template, 3) -CUML_FIL_INITIALIZE_DEVICE(extern template, 4) -CUML_FIL_INITIALIZE_DEVICE(extern template, 5) -CUML_FIL_INITIALIZE_DEVICE(extern template, 6) -CUML_FIL_INITIALIZE_DEVICE(extern template, 7) -CUML_FIL_INITIALIZE_DEVICE(extern template, 8) -CUML_FIL_INITIALIZE_DEVICE(extern template, 9) -CUML_FIL_INITIALIZE_DEVICE(extern template, 10) -CUML_FIL_INITIALIZE_DEVICE(extern template, 11) - -} // namespace device_initialization -} // namespace detail -} // namespace fil - -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/device_initialization/gpu.hpp b/cpp/include/cuml/fil/detail/device_initialization/gpu.hpp deleted file mode 100644 index 99a063d6ad..0000000000 --- a/cpp/include/cuml/fil/detail/device_initialization/gpu.hpp +++ /dev/null @@ -1,32 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include -#include -#include -#include - -#include - -namespace ML { -namespace fil { -namespace detail { -namespace device_initialization { - -/* Non-CUDA header declaration of the GPU specialization for device - * initialization - */ -template -std::enable_if_t, - std::bool_constant>, - void> -initialize_device(raft_proto::device_id device); - -} // namespace device_initialization -} // namespace detail -} // namespace fil - -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/evaluate_tree.hpp b/cpp/include/cuml/fil/detail/evaluate_tree.hpp deleted file mode 100644 index e35d31ad7b..0000000000 --- a/cpp/include/cuml/fil/detail/evaluate_tree.hpp +++ /dev/null @@ -1,201 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -#include -#ifndef __CUDACC__ -#include -#endif -#include -#include -namespace ML { -namespace fil { -namespace detail { - -/* - * Evaluate a single tree on a single row. - * If node_id_mapping is not-nullptr, this kernel outputs leaf node's ID - * instead of the leaf value. - * - * @tparam has_vector_leaves Whether or not this tree has vector leaves - * @tparam has_categorical_nodes Whether or not this tree has any nodes with - * categorical splits - * @tparam node_t The type of nodes in this tree - * @tparam io_t The type used for input to and output from this tree (typically - * either floats or doubles) - * @tparam node_id_mapping_t If non-nullptr_t, this indicates the type we expect for - * node_id_mapping. - * @param node Pointer to the root node of this tree - * @param row Pointer to the input data for this row - * @param first_root_node Pointer to the root node of the first tree. - * @param node_id_mapping Array representing the mapping from internal node IDs to - * final leaf ID outputs - */ -template -HOST DEVICE auto evaluate_tree_impl(node_t const* __restrict__ node, - io_t const* __restrict__ row, - node_t const* __restrict__ first_root_node = nullptr, - node_id_mapping_t node_id_mapping = nullptr) -{ - using categorical_set_type = bitset; - auto cur_node = *node; - do { - auto input_val = row[cur_node.feature_index()]; - auto condition = true; - if constexpr (has_categorical_nodes) { - if (cur_node.is_categorical()) { - auto valid_categories = categorical_set_type{ - &cur_node.index(), uint32_t(sizeof(typename node_t::index_type) * 8)}; - condition = valid_categories.test(input_val) && !isnan(input_val); - } else { - condition = (input_val < cur_node.threshold()); - } - } else { - condition = (input_val < cur_node.threshold()); - } - if (!condition && cur_node.default_distant()) { condition = isnan(input_val); } - node += cur_node.child_offset(condition); - cur_node = *node; - } while (!cur_node.is_leaf()); - if constexpr (std::is_same_v) { - return cur_node.template output(); - } else { - return node_id_mapping[node - first_root_node]; - } -} - -/* - * Evaluate a single tree which requires external categorical storage on a - * single node. - * If node_id_mapping is not-nullptr, this kernel outputs leaf node's ID - * instead of the leaf value. - * - * For non-categorical models and models with a relatively small number of - * categories for any feature, all information necessary for model evaluation - * can be stored on a single node. If the number of categories for any - * feature exceeds the available space on a node, however, the - * categorical split data must be stored external to the node. We pass a - * pointer to this external data and reconstruct bitsets from it indicating - * the positive and negative categories for each categorical node. - * - * @tparam has_vector_leaves Whether or not this tree has vector leaves - * @tparam node_t The type of nodes in this tree - * @tparam io_t The type used for input to and output from this tree (typically - * either floats or doubles) - * @tparam categorical_storage_t The underlying type used for storing - * categorical data (typically char) - * @tparam node_id_mapping_t If non-nullptr_t, this indicates the type we expect for - * node_id_mapping. - * @param node Pointer to the root node of this tree - * @param row Pointer to the input data for this row - * @param categorical_storage Pointer to where categorical split data is - * stored. - */ -template -HOST DEVICE auto evaluate_tree_impl(node_t const* __restrict__ node, - io_t const* __restrict__ row, - categorical_storage_t const* __restrict__ categorical_storage, - node_t const* __restrict__ first_root_node = nullptr, - node_id_mapping_t node_id_mapping = nullptr) -{ - using categorical_set_type = bitset; - auto cur_node = *node; - do { - auto input_val = row[cur_node.feature_index()]; - auto condition = cur_node.default_distant(); - if (!isnan(input_val)) { - if (cur_node.is_categorical()) { - auto valid_categories = - categorical_set_type{categorical_storage + cur_node.index() + 1, - uint32_t(categorical_storage[cur_node.index()])}; - condition = valid_categories.test(input_val); - } else { - condition = (input_val < cur_node.threshold()); - } - } - node += cur_node.child_offset(condition); - cur_node = *node; - } while (!cur_node.is_leaf()); - if constexpr (std::is_same_v) { - return cur_node.template output(); - } else { - return node_id_mapping[node - first_root_node]; - } -} - -/** - * Dispatch to an appropriate version of evaluate_tree kernel. - * - * @tparam has_vector_leaves Whether or not this tree has vector leaves - * @tparam has_categorical_nodes Whether or not this tree has any nodes with - * categorical splits - * @tparam has_nonlocal_categories Whether or not this tree has any nodes that store - * categorical split data externally - * @tparam predict_leaf Whether to predict leaf IDs - * @tparam forest_t The type of forest - * @tparam io_t The type used for input to and output from this tree (typically - * either floats or doubles) - * @tparam categorical_data_t The type for non-local categorical data storage. - * @param forest The forest used to perform inference - * @param tree_index The index of the tree we are evaluating - * @param row The data row we are evaluating - * @param categorical_data The pointer to where non-local data on categorical splits are stored. - */ -template -HOST DEVICE auto evaluate_tree(forest_t const& forest, - index_type tree_index, - io_t const* __restrict__ row, - categorical_data_t categorical_data) -{ - using node_t = typename forest_t::node_type; - if constexpr (predict_leaf) { - auto leaf_node_id = index_type{}; - if constexpr (has_nonlocal_categories) { - leaf_node_id = evaluate_tree_impl(forest.get_tree_root(tree_index), - row, - categorical_data, - forest.get_tree_root(0), - forest.get_node_id_mapping()); - } else { - leaf_node_id = evaluate_tree_impl( - forest.get_tree_root(tree_index), - row, - forest.get_tree_root(0), - forest.get_node_id_mapping()); - } - return leaf_node_id; - } else { - auto tree_output = std::conditional_t{}; - if constexpr (has_nonlocal_categories) { - tree_output = evaluate_tree_impl( - forest.get_tree_root(tree_index), row, categorical_data); - } else { - tree_output = evaluate_tree_impl( - forest.get_tree_root(tree_index), row); - } - return tree_output; - } -} - -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/forest.hpp b/cpp/include/cuml/fil/detail/forest.hpp deleted file mode 100644 index 40f3cea03f..0000000000 --- a/cpp/include/cuml/fil/detail/forest.hpp +++ /dev/null @@ -1,77 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include - -#include - -#include - -namespace ML { -namespace fil { - -/* A collection of trees which together form a forest model - */ -template -struct forest { - using node_type = node; - using io_type = threshold_t; - template - using raw_output_type = std::conditional_t, - std::remove_pointer_t, - typename node_type::threshold_type>; - - HOST DEVICE forest(node_type* forest_nodes, - index_type* forest_root_indexes, - index_type* node_id_mapping, - io_type* bias, - index_type num_trees, - index_type num_outputs) - : nodes_{forest_nodes}, - root_node_indexes_{forest_root_indexes}, - node_id_mapping_{node_id_mapping}, - bias_{bias}, - num_trees_{num_trees}, - num_outputs_{num_outputs} - { - } - - /* Return pointer to the root node of the indicated tree */ - HOST DEVICE auto* get_tree_root(index_type tree_index) const - { - return nodes_ + root_node_indexes_[tree_index]; - } - - /* Return pointer to the mapping from internal node IDs to final node ID outputs. - * Only used when infer_type == infer_kind::leaf_id */ - HOST DEVICE const auto* get_node_id_mapping() const { return node_id_mapping_; } - - /* Return pointer to the bias term */ - HOST DEVICE const auto* bias() const { return bias_; } - - /* Return the number of trees in this forest */ - HOST DEVICE auto tree_count() const { return num_trees_; } - - /* Return the number of outputs per row for default evaluation of this - * forest */ - HOST DEVICE auto num_outputs() const { return num_outputs_; } - - private: - node_type* nodes_; - index_type* root_node_indexes_; - index_type* node_id_mapping_; - io_type* bias_; - index_type num_trees_; - index_type num_outputs_; -}; - -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/gpu_introspection.hpp b/cpp/include/cuml/fil/detail/gpu_introspection.hpp deleted file mode 100644 index dc03e7154f..0000000000 --- a/cpp/include/cuml/fil/detail/gpu_introspection.hpp +++ /dev/null @@ -1,108 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include - -#include - -#include - -namespace ML { -namespace fil { -namespace detail { - -inline auto get_max_shared_mem_per_block( - raft_proto::device_id device_id) -{ - auto thread_local cache = std::vector{}; - if (cache.size() == 0) { - auto device_count = int{}; - raft_proto::cuda_check(cudaGetDeviceCount(&device_count)); - cache.resize(device_count); - for (auto dev = 0; dev < device_count; ++dev) { - raft_proto::cuda_check( - cudaDeviceGetAttribute(&(cache[dev]), cudaDevAttrMaxSharedMemoryPerBlockOptin, dev)); - } - } - return index_type(cache.at(device_id.value())); -} - -inline auto get_sm_count(raft_proto::device_id device_id) -{ - auto thread_local cache = std::vector{}; - if (cache.size() == 0) { - auto device_count = int{}; - raft_proto::cuda_check(cudaGetDeviceCount(&device_count)); - cache.resize(device_count); - for (auto dev = 0; dev < device_count; ++dev) { - raft_proto::cuda_check( - cudaDeviceGetAttribute(&(cache[dev]), cudaDevAttrMultiProcessorCount, dev)); - } - } - return index_type(cache.at(device_id.value())); -} - -inline auto get_max_threads_per_sm(raft_proto::device_id device_id) -{ - auto result = int{}; - raft_proto::cuda_check( - cudaDeviceGetAttribute(&result, cudaDevAttrMaxThreadsPerMultiProcessor, device_id.value())); - return index_type(result); -} - -inline auto get_max_shared_mem_per_sm(raft_proto::device_id device_id) -{ - auto thread_local cache = std::vector{}; - if (cache.size() == 0) { - auto device_count = int{}; - raft_proto::cuda_check(cudaGetDeviceCount(&device_count)); - cache.resize(device_count); - for (auto dev = 0; dev < device_count; ++dev) { - raft_proto::cuda_check( - cudaDeviceGetAttribute(&(cache[dev]), cudaDevAttrMaxSharedMemoryPerMultiprocessor, dev)); - } - } - return index_type(cache.at(device_id.value())); -} - -inline auto get_mem_clock_rate(raft_proto::device_id device_id) -{ - auto result = int{}; - raft_proto::cuda_check( - cudaDeviceGetAttribute(&result, cudaDevAttrMemoryClockRate, device_id.value())); - return index_type(result); -} - -inline auto get_core_clock_rate(raft_proto::device_id device_id) -{ - auto result = int{}; - raft_proto::cuda_check(cudaDeviceGetAttribute(&result, cudaDevAttrClockRate, device_id.value())); - return index_type(result); -} - -/* The maximum number of bytes that can be read in a single instruction */ -auto constexpr static const MAX_READ_CHUNK = index_type{128}; -auto constexpr static const MAX_BLOCKS = index_type{65536}; -auto constexpr static const WARP_SIZE = index_type{32}; -auto constexpr static const MAX_THREADS_PER_BLOCK = index_type{256}; -#ifdef __CUDACC__ -#if __CUDA_ARCH__ == 720 || __CUDA_ARCH__ == 750 || __CUDA_ARCH__ == 860 || \ - __CUDA_ARCH__ == 870 || __CUDA_ARCH__ == 890 || __CUDA_ARCH__ == 1200 || __CUDA_ARCH__ == 1210 -auto constexpr static const MAX_THREADS_PER_SM = index_type{1024}; -#else -auto constexpr static const MAX_THREADS_PER_SM = index_type{2048}; -#endif -#else -auto constexpr static const MAX_THREADS_PER_SM = index_type{2048}; -#endif - -auto constexpr static const MIN_BLOCKS_PER_SM = MAX_THREADS_PER_SM / MAX_THREADS_PER_BLOCK; - -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/index_type.hpp b/cpp/include/cuml/fil/detail/index_type.hpp deleted file mode 100644 index 73dd735291..0000000000 --- a/cpp/include/cuml/fil/detail/index_type.hpp +++ /dev/null @@ -1,11 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -namespace ML { -namespace fil { -using index_type = uint32_t; -} -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/infer.hpp b/cpp/include/cuml/fil/detail/infer.hpp deleted file mode 100644 index f175e6edac..0000000000 --- a/cpp/include/cuml/fil/detail/infer.hpp +++ /dev/null @@ -1,169 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include - -#ifdef CUML_ENABLE_GPU -#include -#endif - -namespace ML { -namespace fil { -namespace detail { - -/* - * Perform inference based on the given forest and input parameters - * - * @tparam D The device type (CPU/GPU) used to perform inference - * @tparam forest_t The type of the forest - * @param forest The forest to be evaluated - * @param postproc The postprocessor object used to execute - * postprocessing - * @param output Pointer to where the output should be written - * @param input Pointer to where the input data can be read from - * @param row_count The number of rows in the input data - * @param col_count The number of columns in the input data - * @param output_count The number of outputs per row - * @param has_categorical_nodes Whether or not any node within the forest has - * a categorical split - * @param vector_output Pointer to the beginning of storage for vector - * outputs of leaves (nullptr for no vector output) - * @param categorical_data Pointer to external categorical data storage if - * required - * @param infer_type Type of inference to perform. Defaults to summing the outputs of all trees - * and produce an output per row. If set to "per_tree", we will instead output all outputs of - * individual trees. If set to "leaf_id", we will output the integer ID of the leaf node - * for each tree. - * @param specified_chunk_size If non-nullopt, the size of "mini-batches" - * used for distributing work across threads - * @param device The device on which to execute evaluation - * @param stream Optionally, the CUDA stream to use - */ -template -void infer(forest_t const& forest, - postprocessor const& postproc, - typename forest_t::io_type* output, - typename forest_t::io_type* input, - index_type row_count, - index_type col_count, - index_type output_count, - bool has_categorical_nodes, - typename forest_t::io_type* vector_output = nullptr, - typename forest_t::node_type::index_type* categorical_data = nullptr, - infer_kind infer_type = infer_kind::default_kind, - std::optional specified_chunk_size = std::nullopt, - raft_proto::device_id device = raft_proto::device_id{}, - raft_proto::cuda_stream stream = raft_proto::cuda_stream{}) -{ - if (vector_output == nullptr) { - if (categorical_data == nullptr) { - if (!has_categorical_nodes) { - inference::infer(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - nullptr, - nullptr, - infer_type, - specified_chunk_size, - device, - stream); - } else { - inference::infer(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - nullptr, - nullptr, - infer_type, - specified_chunk_size, - device, - stream); - } - } else { - inference::infer(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - nullptr, - categorical_data, - infer_type, - specified_chunk_size, - device, - stream); - } - } else { - if (categorical_data == nullptr) { - if (!has_categorical_nodes) { - inference::infer(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - vector_output, - nullptr, - infer_type, - specified_chunk_size, - device, - stream); - } else { - inference::infer(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - vector_output, - nullptr, - infer_type, - specified_chunk_size, - device, - stream); - } - } else { - inference::infer(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - vector_output, - categorical_data, - infer_type, - specified_chunk_size, - device, - stream); - } - } -} - -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/infer/cpu.hpp b/cpp/include/cuml/fil/detail/infer/cpu.hpp deleted file mode 100644 index 344ea60c1b..0000000000 --- a/cpp/include/cuml/fil/detail/infer/cpu.hpp +++ /dev/null @@ -1,147 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include - -namespace ML { -namespace fil { -namespace detail { -namespace inference { - -/* A wrapper around the underlying inference kernels to support dispatching to - * the right kernel - * - * This specialization is used for CPU inference and for requests for GPU - * inference on non-GPU-enabled builds. An exception will be thrown if a - * request is made for GPU on inference on a non-GPU-enabled build. - * - * @tparam D The type of device (CPU/GPU) on which to perform inference. - * @tparam has_categorical_nodes Whether or not any node in the model has - * categorical splits. - * @tparam vector_output_t If non-nullptr_t, the type of vector leaf output - * @tparam categorical_data_t If non-nullptr_t, the type of non-local - * categorical data storage - * - * @param forest The forest to be used for inference. - * @param postproc The postprocessor object to be used for postprocessing raw - * output from the forest. - * @param row_count The number of rows in the input - * @param col_count The number of columns per row in the input - * @param output_count The number of output elements per row - * @param vector_output If non-nullptr, a pointer to storage for vector leaf - * outputs - * @param categorical_data If non-nullptr, a pointer to non-local storage for - * data on categorical splits. - * @param infer_type Type of inference to perform. Defaults to summing the outputs of all trees - * and produce an output per row. If set to "per_tree", we will instead output all outputs of - * individual trees. If set to "leaf_id", we will output the integer ID of the leaf node - * for each tree. - * @param specified_chunk_size If non-nullopt, the mini-batch size used for - * processing rows in a batch. For CPU inference, this essentially determines - * the granularity of parallelism. A larger chunk size means that a single - * thread will process more rows for its assigned trees before fetching a - * new batch of rows. In general, so long as the chunk size remains much - * smaller than the batch size (minimally less than the batch size divided by - * the number of available cores), larger batches see improved performance with - * larger chunk sizes. Unlike for GPU, any positive value is valid (up to - * hardware constraints), but it is recommended to test powers of 2 from 1 - * (for individual row inference) to 512 (for very large batch - * inference). A value of 64 is a generally-useful default. - */ -template -std::enable_if_t, - std::bool_constant>, - void> -infer(forest_t const& forest, - postprocessor const& postproc, - typename forest_t::io_type* output, - typename forest_t::io_type* input, - index_type row_count, - index_type col_count, - index_type output_count, - vector_output_t vector_output = nullptr, - categorical_data_t categorical_data = nullptr, - infer_kind infer_type = infer_kind::default_kind, - std::optional specified_chunk_size = std::nullopt, - raft_proto::device_id device = raft_proto::device_id{}, - raft_proto::cuda_stream = raft_proto::cuda_stream{}) -{ - if constexpr (D == raft_proto::device_type::gpu) { - throw raft_proto::gpu_unsupported("Tried to use GPU inference in CPU-only build"); - } else { - if (infer_type == infer_kind::leaf_id) { - infer_kernel_cpu( - forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - specified_chunk_size.value_or(hardware_constructive_interference_size), - hardware_constructive_interference_size, - vector_output, - categorical_data, - infer_type); - } else { - infer_kernel_cpu( - forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - specified_chunk_size.value_or(hardware_constructive_interference_size), - hardware_constructive_interference_size, - vector_output, - categorical_data, - infer_type); - } - } -} - -/* This macro is invoked here to declare all standard specializations of this - * template as extern. This ensures that this (relatively complex) code is - * compiled as few times as possible. A macro is used because ever - * specialization must be explicitly declared. The final argument to the macro - * references the 8 specialization variants compiled in standard cuML FIL. */ -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 0) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 1) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 2) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 3) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 4) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 5) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 6) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 7) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 8) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 9) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 10) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::cpu, 11) - -} // namespace inference -} // namespace detail -} // namespace fil - -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/infer/gpu.cuh b/cpp/include/cuml/fil/detail/infer/gpu.cuh deleted file mode 100644 index aa4bbd391f..0000000000 --- a/cpp/include/cuml/fil/detail/infer/gpu.cuh +++ /dev/null @@ -1,345 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include -#include -#include -#include -#include - -namespace ML { -namespace fil { -namespace detail { -namespace inference { - -inline auto compute_output_size(index_type row_output_size, - index_type threads_per_block, - index_type rows_per_block_iteration, - infer_kind infer_type = infer_kind::default_kind) -{ - auto result = row_output_size * rows_per_block_iteration; - if (infer_type == infer_kind::default_kind) { - result *= raft_proto::ceildiv(threads_per_block, rows_per_block_iteration); - } - return result; -} - -/* A wrapper around the underlying inference kernels to support dispatching to - * the right kernel - * - * This specialization is used for GPU inference. It performs any necessary - * computation necessary prior to kernel launch and then launches the correct - * inference kernel. - * - * @tparam D The type of device (CPU/GPU) on which to perform inference. - * @tparam has_categorical_nodes Whether or not any node in the model has - * categorical splits. - * @tparam vector_output_t If non-nullptr_t, the type of vector leaf output - * @tparam categorical_data_t If non-nullptr_t, the type of non-local - * categorical data storage - * - * @param forest The forest to be used for inference. - * @param postproc The postprocessor object to be used for postprocessing raw - * output from the forest. - * @param row_count The number of rows in the input - * @param col_count The number of columns per row in the input - * @param output_count The number of output elements per row - * @param vector_output If non-nullptr, a pointer to storage for vector leaf - * outputs - * @param categorical_data If non-nullptr, a pointer to non-local storage for - * data on categorical splits. - * @param infer_type Type of inference to perform. Defaults to summing the outputs of all trees - * and produce an output per row. If set to "per_tree", we will instead output all outputs of - * individual trees. If set to "leaf_id", we will output the integer ID of the leaf node - * for each tree. - * @param specified_chunk_size If non-nullopt, the mini-batch size used for - * processing rows in a batch. For GPU inference, this determines the number of - * rows that are processed per iteration of inference in a single block. It - * is difficult to predict the optimal value for this parameter, but tuning it - * can result in a substantial improvement in performance. The optimal - * value depends on hardware, model, and batch size. Valid values are any power - * of 2 from 1 to 32. - */ -template -std::enable_if_t infer( - forest_t const& forest, - postprocessor const& postproc, - typename forest_t::io_type* output, - typename forest_t::io_type* input, - index_type row_count, - index_type col_count, - index_type output_count, - vector_output_t vector_output = nullptr, - categorical_data_t categorical_data = nullptr, - infer_kind infer_type = infer_kind::default_kind, - std::optional specified_chunk_size = std::nullopt, - raft_proto::device_id device = raft_proto::device_id{}, - raft_proto::cuda_stream stream = raft_proto::cuda_stream{}) -{ - using output_t = typename forest_t::template raw_output_type; - - auto sm_count = get_sm_count(device); - auto const max_shared_mem_per_block = get_max_shared_mem_per_block(device); - auto const max_shared_mem_per_sm = get_max_shared_mem_per_sm(device); - auto const max_overall_shared_mem = std::min(max_shared_mem_per_block, max_shared_mem_per_sm); - - auto row_size_bytes = index_type(index_type(sizeof(typename forest_t::io_type) * col_count)); - auto row_output_size = output_count; - auto row_output_size_bytes = index_type(sizeof(typename forest_t::io_type) * row_output_size); - - // First determine the number of threads per block. This is the indicated - // preferred value unless we cannot handle at least 1 row per block iteration - // with available shared memory, in which case we must reduce the threads per - // block. - auto threads_per_block = - min(MAX_THREADS_PER_BLOCK, - raft_proto::downpadded_size( - (max_shared_mem_per_block - row_size_bytes) / row_output_size_bytes, WARP_SIZE)); - - // If we cannot do at least a warp per block when storing input rows in - // shared mem, recalculate our threads per block without input storage - if (threads_per_block < WARP_SIZE) { - threads_per_block = - min(MAX_THREADS_PER_BLOCK, - raft_proto::downpadded_size(max_shared_mem_per_block / row_output_size_bytes, WARP_SIZE)); - if (threads_per_block >= WARP_SIZE) { - row_size_bytes = index_type{}; // Do not store input rows in shared mem - } - } - - // If we cannot do at least a warp per block when storing output in - // shared mem, recalculate our threads per block with ONLY input storage - if (threads_per_block < WARP_SIZE) { - threads_per_block = - min(MAX_THREADS_PER_BLOCK, - raft_proto::downpadded_size(max_shared_mem_per_block / row_size_bytes, WARP_SIZE)); - } - - // If we still cannot use at least a warp per block, give up on using - // shared memory and just maximize occupancy - if (threads_per_block < WARP_SIZE) { threads_per_block = MAX_THREADS_PER_BLOCK; } - - auto const max_resident_blocks = sm_count * (get_max_threads_per_sm(device) / threads_per_block); - - // Compute shared memory usage based on minimum or specified - // rows_per_block_iteration - auto rows_per_block_iteration = specified_chunk_size.value_or(index_type{1}); - auto constexpr const output_item_bytes = index_type(sizeof(output_t)); - auto output_workspace_size = - compute_output_size(row_output_size, threads_per_block, rows_per_block_iteration, infer_type); - auto output_workspace_size_bytes = output_item_bytes * output_workspace_size; - auto global_workspace = raft_proto::buffer{}; - - if (output_workspace_size_bytes > max_shared_mem_per_block) { - output_workspace_size_bytes = 0; - row_output_size = 0; - } - auto shared_mem_per_block = - min(rows_per_block_iteration * row_size_bytes + output_workspace_size_bytes, - max_overall_shared_mem); - - auto resident_blocks_per_sm = - min(raft_proto::ceildiv(max_shared_mem_per_sm, shared_mem_per_block), max_resident_blocks); - - // If caller has not specified the number of rows per block iteration, apply - // the following heuristic to identify an approximately optimal value - if (!specified_chunk_size.has_value() && resident_blocks_per_sm >= MIN_BLOCKS_PER_SM) { - rows_per_block_iteration = index_type{32}; - } - - if (row_output_size != 0 && rows_per_block_iteration > 1) { - do { - output_workspace_size = compute_output_size( - row_output_size, threads_per_block, rows_per_block_iteration, infer_type); - output_workspace_size_bytes = output_item_bytes * output_workspace_size; - - shared_mem_per_block = - (rows_per_block_iteration * row_size_bytes + output_workspace_size_bytes); - if (shared_mem_per_block > max_overall_shared_mem) { - rows_per_block_iteration >>= index_type{1}; - } - } while (shared_mem_per_block > max_overall_shared_mem && rows_per_block_iteration > 1); - } - - shared_mem_per_block = std::min(shared_mem_per_block, max_overall_shared_mem); - - // Divide shared mem evenly - shared_mem_per_block = std::min( - max_overall_shared_mem, max_shared_mem_per_sm / (max_shared_mem_per_sm / shared_mem_per_block)); - - auto num_blocks = std::min(raft_proto::ceildiv(row_count, rows_per_block_iteration), MAX_BLOCKS); - if (row_output_size == 0) { - global_workspace = raft_proto::buffer{ - output_workspace_size * num_blocks, raft_proto::device_type::gpu, device.value(), stream}; - } - - /** - * Throw an error for large inputs that would cause integer overflow. - * TODO(hcho3): Support large inputs via streaming - **/ - { - // Use 64-bit integers for intermediate computations, to avoid overflows - // while computing max_num_row. - auto chunk_size = std::uint64_t{32}; - while (rows_per_block_iteration <= chunk_size / 2 && chunk_size >= 2) { - chunk_size /= 2; - } - auto task_count = chunk_size * forest.tree_count(); - auto num_grove = [infer_type, threads_per_block, task_count, chunk_size]() { - auto result = std::uint64_t{1}; - if (infer_type == infer_kind::default_kind) { - result = raft_proto::ceildiv(min(static_cast(threads_per_block), task_count), - chunk_size); - } - return result; - }(); - auto max_num_row = static_cast(std::numeric_limits::max()) / - (output_count * num_grove); - if (max_num_row >= 3) { - max_num_row -= 3; - // -3 is part of the upper bound on num_row, to ensure that the offset - // does not overflow past the uint32_t limit. - } - if (row_count > max_num_row) { - throw runtime_error(std::string("Input size too large! Input should be at most ") + - std::to_string(max_num_row) + "."); - } - } - - if (rows_per_block_iteration <= 1) { - infer_kernel - <<>>(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - shared_mem_per_block, - output_workspace_size, - vector_output, - categorical_data, - infer_type, - global_workspace.data()); - } else if (rows_per_block_iteration <= 2) { - infer_kernel - <<>>(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - shared_mem_per_block, - output_workspace_size, - vector_output, - categorical_data, - infer_type, - global_workspace.data()); - } else if (rows_per_block_iteration <= 4) { - infer_kernel - <<>>(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - shared_mem_per_block, - output_workspace_size, - vector_output, - categorical_data, - infer_type, - global_workspace.data()); - } else if (rows_per_block_iteration <= 8) { - infer_kernel - <<>>(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - shared_mem_per_block, - output_workspace_size, - vector_output, - categorical_data, - infer_type, - global_workspace.data()); - } else if (rows_per_block_iteration <= 16) { - infer_kernel - <<>>(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - shared_mem_per_block, - output_workspace_size, - vector_output, - categorical_data, - infer_type, - global_workspace.data()); - } else { - infer_kernel - <<>>(forest, - postproc, - output, - input, - row_count, - col_count, - output_count, - shared_mem_per_block, - output_workspace_size, - vector_output, - categorical_data, - infer_type, - global_workspace.data()); - } - raft_proto::cuda_check(cudaGetLastError()); -} - -/* This macro is invoked here to declare all standard specializations of this - * template as extern. This ensures that this (relatively complex) code is - * compiled as few times as possible. A macro is used because ever - * specialization must be explicitly declared. The final argument to the macro - * references the 8 specialization variants compiled in standard cuML FIL. */ -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::gpu, 0) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::gpu, 1) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::gpu, 2) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::gpu, 3) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::gpu, 4) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::gpu, 5) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::gpu, 6) -CUML_FIL_INFER_ALL(extern template, raft_proto::device_type::gpu, 7) - -} // namespace inference -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/infer/gpu.hpp b/cpp/include/cuml/fil/detail/infer/gpu.hpp deleted file mode 100644 index 6260c44127..0000000000 --- a/cpp/include/cuml/fil/detail/infer/gpu.hpp +++ /dev/null @@ -1,45 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -#include -#include -#include - -#include -#include - -namespace ML { -namespace fil { -namespace detail { -namespace inference { - -/* The CUDA-free header declaration of the GPU infer template */ -template -std::enable_if_t infer( - forest_t const& forest, - postprocessor const& postproc, - typename forest_t::io_type* output, - typename forest_t::io_type* input, - index_type row_count, - index_type col_count, - index_type class_count, - vector_output_t vector_output = nullptr, - categorical_data_t categorical_data = nullptr, - infer_kind infer_type = infer_kind::default_kind, - std::optional specified_chunk_size = std::nullopt, - raft_proto::device_id device = raft_proto::device_id{}, - raft_proto::cuda_stream stream = raft_proto::cuda_stream{}); - -} // namespace inference -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/infer_kernel/cpu.hpp b/cpp/include/cuml/fil/detail/infer_kernel/cpu.hpp deleted file mode 100644 index 3825ddb41f..0000000000 --- a/cpp/include/cuml/fil/detail/infer_kernel/cpu.hpp +++ /dev/null @@ -1,203 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include -#include -#include -#include - -#include - -#ifdef _OPENMP -#include -#else -#ifdef omp_get_max_threads -#if omp_get_max_threads() != 1 -#error "Inconsistent placeholders for omp_get_max_threads" -#endif -#else -#define omp_get_max_threads() 1 -#endif -#endif - -#include -#include -#include -#include -#include -#include -#include -#include - -namespace ML { -namespace fil { -namespace detail { - -/** - * The CPU "kernel" used to actually perform forest inference - * - * @tparam has_categorical_nodes Whether or not this kernel should be - * compiled to operate on trees with categorical nodes. - * @tparam forest_t The type of the forest object which will be used for - * inference. - * @tparam vector_output_t If non-nullptr_t, this indicates the type we expect - * for outputs from vector leaves. - * @tparam categorical_data_t If non-nullptr_t, this indicates the type we - * expect for non-local categorical data storage. - * @param forest The forest used to perform inference - * @param postproc The postprocessor object used to store all necessary - * data for postprocessing - * @param output Pointer to the host-accessible buffer where output - * should be written - * @param input Pointer to the host-accessible buffer where input should be - * read from - * @param row_count The number of rows in the input - * @param col_count The number of columns per row in the input - * @param num_outputs The expected number of output elements per row - * @param chunk_size The number of rows for each thread to process with its - * assigned trees before fetching a new set of trees/rows. - * @param grove_size The number of trees to assign to a thread for each chunk - * of rows it processes. - * @param vector_output_p If non-nullptr, a pointer to the stored leaf - * vector outputs for all leaf nodes - * @param categorical_data If non-nullptr, a pointer to where non-local - * data on categorical splits are stored. - * @param infer_type Type of inference to perform. Defaults to summing the outputs of all trees - * and produce an output per row. If set to "per_tree", we will instead output all outputs of - * individual trees. If set to "leaf_id", we will output the integer ID of the leaf node - * for each tree. - */ -template -void infer_kernel_cpu(forest_t const& forest, - postprocessor const& postproc, - typename forest_t::io_type* output, - typename forest_t::io_type const* input, - index_type row_count, - index_type col_count, - index_type num_outputs, - index_type chunk_size = hardware_constructive_interference_size, - index_type grove_size = hardware_constructive_interference_size, - vector_output_t vector_output_p = nullptr, - categorical_data_t categorical_data = nullptr, - infer_kind infer_type = infer_kind::default_kind) -{ - auto constexpr has_vector_leaves = !std::is_same_v; - auto constexpr has_nonlocal_categories = !std::is_same_v; - - using node_t = typename forest_t::node_type; - - using output_t = typename forest_t::template raw_output_type; - - auto const num_tree = forest.tree_count(); - auto const num_grove = raft_proto::ceildiv(num_tree, grove_size); - auto const num_chunk = raft_proto::ceildiv(row_count, chunk_size); - - /** - * Throw an error for large inputs that would cause integer overflow. - * TODO(hcho3): Support large inputs via streaming - **/ - { - // Use 64-bit integers for intermediate computations, to avoid overflows - // while computing max_num_row. - auto max_num_row = static_cast(std::numeric_limits::max()) / - (num_outputs * static_cast(num_grove)); - if (max_num_row >= 3) { - max_num_row -= 3; - // -3 is part of the upper bound on num_row, to ensure that the offset - // does not overflow past the uint32_t limit. - } - if (row_count > max_num_row) { - throw runtime_error(std::string("Input size too large! Input should be at most ") + - std::to_string(max_num_row) + "."); - } - } - - auto output_workspace = std::vector(row_count * num_outputs * num_grove, output_t{}); - auto const task_count = num_grove * num_chunk; - -#pragma omp parallel num_threads(std::min(index_type(omp_get_max_threads()), task_count)) - { - // Infer on each grove and chunk -#pragma omp for - for (auto task_index = index_type{}; task_index < task_count; ++task_index) { - auto const grove_index = task_index / num_chunk; - auto const chunk_index = task_index % num_chunk; - auto const start_row = chunk_index * chunk_size; - auto const end_row = std::min(start_row + chunk_size, row_count); - auto const start_tree = grove_index * grove_size; - auto const end_tree = std::min(start_tree + grove_size, num_tree); - - for (auto row_index = start_row; row_index < end_row; ++row_index) { - for (auto tree_index = start_tree; tree_index < end_tree; ++tree_index) { - auto tree_output = - std::conditional_t>{}; - tree_output = evaluate_tree( - forest, tree_index, input + row_index * col_count, categorical_data); - if constexpr (predict_leaf) { - output_workspace[row_index * num_outputs * num_grove + tree_index * num_grove + - grove_index] = static_cast(tree_output); - } else { - auto const default_num_outputs = forest.num_outputs(); - if constexpr (has_vector_leaves) { - auto output_offset = (row_index * num_outputs * num_grove + - tree_index * default_num_outputs * num_grove * - (infer_type == infer_kind::per_tree) + - grove_index); - for (auto output_index = index_type{}; output_index < default_num_outputs; - ++output_index) { - output_workspace[output_offset + output_index * num_grove] += - vector_output_p[tree_output * default_num_outputs + output_index]; - } - } else { - auto output_offset = - (row_index * num_outputs * num_grove + - (tree_index % default_num_outputs) * num_grove * - (infer_type == infer_kind::default_kind) + - tree_index * num_grove * (infer_type == infer_kind::per_tree) + grove_index); - output_workspace[output_offset] += tree_output; - } - } - } // Trees - } // Rows - } // Tasks - - // Sum over grove and postprocess -#pragma omp for - for (auto row_index = index_type{}; row_index < row_count; ++row_index) { - for (auto output_index = index_type{}; output_index < num_outputs; ++output_index) { - auto grove_offset = (row_index * num_outputs * num_grove + output_index * num_grove); - - output_workspace[grove_offset] = - std::accumulate(std::begin(output_workspace) + grove_offset, - std::begin(output_workspace) + grove_offset + num_grove, - output_t{}); - } - postproc(infer_type, - output_workspace.data() + row_index * num_outputs * num_grove, - num_outputs, - forest.bias(), - output + row_index * num_outputs, - num_grove); - } - } // End omp parallel -} - -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/infer_kernel/gpu.cuh b/cpp/include/cuml/fil/detail/infer_kernel/gpu.cuh deleted file mode 100644 index 95801160da..0000000000 --- a/cpp/include/cuml/fil/detail/infer_kernel/gpu.cuh +++ /dev/null @@ -1,213 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include - -namespace ML { -namespace fil { -namespace detail { - -/** - * The GPU kernel used to actually perform forest inference - * - * @tparam has_categorical_nodes Whether or not this kernel should be - * compiled to operate on trees with categorical nodes. - * @tparam chunk_size The number of rows to be simultaneously processed - * in each iteration of inference within a single block. This is a - * performance tuning parameter, and having it fixed at compile-time offers a - * measurable performance benefit. In standard cuML FIL, we compile for all - * powers of 2 from 1 to 32. A power of 2 is not guaranteed to optimize - * performance for all batch sizes and models, but it is far more likely to - * than other values. - * @tparam forest_t The type of the forest object which will be used for - * inference. - * @tparam vector_output_t If non-nullptr_t, this indicates the type we expect - * for outputs from vector leaves. - * @tparam categorical_data_t If non-nullptr_t, this indicates the type we - * expect for non-local categorical data storage. - * @param forest The forest used to perform inference - * @param postproc The postprocessor object used to store all necessary - * data for postprocessing - * @param output Pointer to the device-accessible buffer where output - * should be written - * @param input Pointer to the device-accessible buffer where input should be - * read from - * @param row_count The number of rows in the input - * @param col_count The number of columns per row in the input - * @param num_outputs The expected number of output elements per row - * @param shared_mem_byte_size The number of bytes of shared memory allocated - * to this kernel. - * @param output_workspace_size The total number of temporary elements required - * to be stored as an intermediate output during inference - * @param vector_output_p If non-nullptr, a pointer to the stored leaf - * vector outputs for all leaf nodes - * @param categorical_data If non-nullptr, a pointer to where non-local - * data on categorical splits are stored. - * @param infer_type Type of inference to perform. Defaults to summing the outputs of all trees - * and produce an output per row. If set to "per_tree", we will instead output all outputs of - * individual trees. If set to "leaf_id", we will instead output the integer ID of the leaf node - * for each tree. - * @param global_mem_fallback_buffer Buffer to use as a fallback, when there isn't enough shared - * memory. Set it to nullptr to disable - */ -template -CUML_KERNEL void __launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_SM) infer_kernel( - forest_t forest, - postprocessor postproc, - typename forest_t::io_type* output, - typename forest_t::io_type const* input, - index_type row_count, - index_type col_count, - index_type num_outputs, - index_type shared_mem_byte_size, - index_type output_workspace_size, - vector_output_t vector_output_p = nullptr, - categorical_data_t categorical_data = nullptr, - infer_kind infer_type = infer_kind::default_kind, - typename forest_t::template raw_output_type* workspace_fallback = nullptr) -{ - auto const default_num_outputs = forest.num_outputs(); - auto constexpr has_vector_leaves = !std::is_same_v; - auto constexpr has_nonlocal_categories = !std::is_same_v; - using output_t = typename forest_t::template raw_output_type; - extern __shared__ std::byte shared_mem_raw[]; - - auto shared_mem = shared_memory_buffer(shared_mem_raw, shared_mem_byte_size); - - using node_t = typename forest_t::node_type; - - using io_t = typename forest_t::io_type; - - for (auto i = blockIdx.x * chunk_size; i < row_count; i += chunk_size * gridDim.x) { - // i: the ID of the first row in the current chunk - - shared_mem.clear(); - auto* output_workspace = shared_mem.fill( - output_workspace_size, output_t{}, (workspace_fallback + blockIdx.x * output_workspace_size)); - - // Handle as many rows as requested per loop or as many rows as are left to - // process - auto rows_in_this_iteration = min(chunk_size, row_count - i); - - auto* input_data = shared_mem.copy(input + i * col_count, rows_in_this_iteration, col_count); - - auto task_count = chunk_size * forest.tree_count(); - - auto num_grove = raft_proto::ceildiv(min(index_type(blockDim.x), task_count), chunk_size) * - (infer_type == infer_kind::default_kind) + - (infer_type != infer_kind::default_kind); - - // Note that this sync is safe because every thread in the block will agree - // on whether or not a sync is required - shared_mem.sync(); - - // Every thread must iterate the same number of times in order to avoid a - // deadlock on __syncthreads, so we round the task_count up to the next - // multiple of the number of threads in this block. We then only perform - // work within the loop if the task_index is below the actual task_count. - auto const task_count_rounded_up = blockDim.x * raft_proto::ceildiv(task_count, blockDim.x); - - // Infer on each tree and row - for (auto task_index = threadIdx.x; task_index < task_count_rounded_up; - task_index += blockDim.x) { - auto row_index = task_index % chunk_size; - auto real_task = task_index < task_count && row_index < rows_in_this_iteration; - row_index *= real_task; - auto tree_index = task_index * real_task / chunk_size; - auto grove_index = (threadIdx.x / chunk_size) * (infer_type == infer_kind::default_kind); - - auto tree_output = std::conditional_t{}; - auto leaf_node_id = index_type{}; - if (infer_type == infer_kind::leaf_id) { - leaf_node_id = - evaluate_tree( - forest, tree_index, input_data + row_index * col_count, categorical_data); - } else { - tree_output = - evaluate_tree( - forest, tree_index, input_data + row_index * col_count, categorical_data); - } - - if (infer_type == infer_kind::leaf_id) { - output_workspace[row_index * num_outputs * num_grove + tree_index * num_grove + - grove_index] = static_cast(leaf_node_id); - } else { - if constexpr (has_vector_leaves) { - auto output_offset = - (row_index * num_outputs * num_grove + - tree_index * default_num_outputs * num_grove * (infer_type == infer_kind::per_tree) + - grove_index); - for (auto output_index = index_type{}; output_index < default_num_outputs; - ++output_index) { - if (real_task) { - output_workspace[output_offset + output_index * num_grove] += - vector_output_p[tree_output * default_num_outputs + output_index]; - } - } - } else { - auto output_offset = - (row_index * num_outputs * num_grove + - (tree_index % default_num_outputs) * num_grove * - (infer_type == infer_kind::default_kind) + - tree_index * num_grove * (infer_type == infer_kind::per_tree) + grove_index); - if (real_task) { output_workspace[output_offset] += tree_output; } - } - } - - __syncthreads(); - } - - auto padded_num_groves = raft_proto::padded_size(num_grove, WARP_SIZE); - for (auto row_index = threadIdx.x / WARP_SIZE; row_index < rows_in_this_iteration; - row_index += blockDim.x / WARP_SIZE) { - for (auto class_index = index_type{}; class_index < num_outputs; ++class_index) { - auto grove_offset = (row_index * num_outputs * num_grove + class_index * num_grove); - auto class_sum = output_t{}; - for (auto grove_index = threadIdx.x % WARP_SIZE; grove_index < padded_num_groves; - grove_index += WARP_SIZE) { - auto real_thread = grove_index < num_grove; - auto out_index = grove_offset + grove_index * real_thread; - class_sum *= (threadIdx.x % WARP_SIZE == 0); - class_sum += output_workspace[out_index] * real_thread; - for (auto thread_offset = (WARP_SIZE >> 1); thread_offset > 0; thread_offset >>= 1) { - class_sum += __shfl_down_sync(0xFFFFFFFF, class_sum, thread_offset); - } - } - if (threadIdx.x % WARP_SIZE == 0) { output_workspace[grove_offset] = class_sum; } - } - if (threadIdx.x % WARP_SIZE == 0) { - postproc(infer_type, - output_workspace + row_index * num_outputs * num_grove, - num_outputs, - forest.bias(), - output + ((i + row_index) * num_outputs), - num_grove); - } - } - __syncthreads(); - } -} - -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/infer_kernel/shared_memory_buffer.cuh b/cpp/include/cuml/fil/detail/infer_kernel/shared_memory_buffer.cuh deleted file mode 100644 index 20f0c3c758..0000000000 --- a/cpp/include/cuml/fil/detail/infer_kernel/shared_memory_buffer.cuh +++ /dev/null @@ -1,147 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -#include - -#include -#include - -namespace ML { -namespace fil { - -/* A struct used to simplify complex access to a buffer of shared memory - * - * @param buffer A pointer to the shared memory allocation - * @param size The size in bytes of the shared memory allocation - */ -struct shared_memory_buffer { - __device__ shared_memory_buffer(std::byte* buffer = nullptr, index_type size = index_type{}) - : data{buffer}, total_size{size}, remaining_data{buffer}, remaining_size{size} - { - } - - /* If possible, copy the given number of rows with the given number of columns from source - * to the end of this buffer, padding each row by the given number of - * elements (usually to reduce memory bank conflicts). If there is not enough - * room, no copy is performed. Return a pointer to the desired data, whether - * that is in the original location or copied to shared memory. */ - template - __device__ auto* copy(T* source, - index_type row_count, - index_type col_count, - index_type row_pad = index_type{}) - { - auto* dest = reinterpret_cast*>(remaining_data); - auto source_count = row_count * col_count; - auto dest_count = row_count * (col_count + row_pad); - - auto copy_data = (dest_count * sizeof(T) <= remaining_size); - - source_count *= copy_data; - for (auto i = threadIdx.x; i < source_count; i += blockDim.x) { - dest[i + row_pad * (i / col_count)] = source[i]; - } - - auto* result = copy_data ? static_cast(dest) : source; - requires_sync = requires_sync || copy_data; - - auto offset = dest_count * index_type(sizeof(T)); - remaining_data += offset; - remaining_size -= offset; - - return result; - } - - /* If possible, copy the given number of elements from source to the end of this buffer - * If there is not enough room, no copy is performed. Return a pointer to the - * desired data, whether that is in the original location or copied to shared - * memory. */ - template - __device__ auto* copy(T* source, index_type element_count) - { - auto* dest = reinterpret_cast*>(remaining_data); - - auto copy_data = (element_count * index_type(sizeof(T)) <= remaining_size); - - element_count *= copy_data; - for (auto i = threadIdx.x; i < element_count; i += blockDim.x) { - dest[i] = source[i]; - } - auto* result = copy_data ? static_cast(dest) : source; - requires_sync = requires_sync || copy_data; - - auto offset = element_count * index_type(sizeof(T)); - remaining_data += offset; - remaining_size -= offset; - - return result; - } - - /* If possible, fill the next element_count elements with given value. If - * there is not enough room, the fill is not performed. Return a pointer to - * the start of the desired data if the fill was possible or else nullptr. */ - template - __device__ auto* fill(index_type element_count, T value = T{}, T* fallback_buffer = nullptr) - { - auto* dest = reinterpret_cast*>(remaining_data); - - auto copy_data = (element_count * index_type(sizeof(T)) <= remaining_size); - - element_count *= copy_data; - for (auto i = threadIdx.x; i < element_count; i += blockDim.x) { - dest[i] = value; - } - - auto* result = copy_data ? static_cast(dest) : fallback_buffer; - requires_sync = requires_sync || copy_data; - - auto offset = element_count * index_type(sizeof(T)); - remaining_data += offset; - remaining_size -= offset; - - return result; - } - - /* Clear all stored data and return a pointer to the beginning of available - * shared memory */ - __device__ auto* clear() - { - remaining_size = total_size; - remaining_data = data; - return remaining_data; - } - - /* Pad stored data to ensure correct alignment for given type */ - template - __device__ void align() - { - auto pad_required = (total_size - remaining_size) % index_type(sizeof(T)); - remaining_data += pad_required; - remaining_size -= pad_required; - } - - /* If necessary, sync threads. Note that this can cause a deadlock if not all - * threads call this method. */ - __device__ void sync() - { - if (requires_sync) { __syncthreads(); } - requires_sync = false; - } - - /* Return the remaining size in bytes left in this buffer */ - __device__ auto remaining() { return remaining_size; } - - private: - std::byte* data; - index_type total_size; - std::byte* remaining_data; - index_type remaining_size; - bool requires_sync; -}; - -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/node.hpp b/cpp/include/cuml/fil/detail/node.hpp deleted file mode 100644 index 0a5bf88e2c..0000000000 --- a/cpp/include/cuml/fil/detail/node.hpp +++ /dev/null @@ -1,248 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include -#include -#include -#include - -#include -#include - -namespace ML { -namespace fil { - -namespace detail { - -/* - * Return the byte size to which a node with the given types should be aligned - */ -template -auto constexpr get_node_alignment() -{ - auto total = index_type(std::max(sizeof(threshold_t), sizeof(index_t)) + - sizeof(metadata_storage_t) + sizeof(offset_t)); - auto result = index_type{8}; - if (total > result) { result = index_type{16}; } - if (total > result) { result = index_type{32}; } - if (total > result) { result = index_type{64}; } - if (total > result) { result = index_type{128}; } - if (total > result) { result = index_type{256}; } - if (total > result) { result = total; } - return result; -} - -} // namespace detail - -/* @brief A single node in a forest model - * - * Note that this implementation includes NO error checking for poorly-chosen - * template types. If the types are not large enough to hold the required data, - * an incorrect node will be silently constructed. Error checking occurs - * instead at the time of construction of the entire forest. - * - * @tparam layout_v The layout for nodes within the forest - * - * @tparam threshold_t The type used as a threshold for evaluating a non-leaf - * node or (when possible) the output of a leaf node. For non-categorical - * nodes, if an input value is less than this threshold, the node evaluates to - * true. For leaf nodes, output values will only be stored as this type if it - * matches the leaf output type expected by the forest. Typically, this type is - * either float or double. - * - * @tparam index_t The type used as an index to the output data for leaf nodes, - * or to the categorical set for a categorical non-leaf node. This type should - * be large enough to index the entire array of output data or categorical sets - * stored in the forest. Typically, this type is either uint32_t or uint64_t. - * Smaller types offer no benefit, since this value is stored in a union with - * threshold_t, which is at least 32 bits. - * - * @tparam metadata_storage_t An unsigned integral type used for a bit-wise - * representation of metadata about this node. The first three bits encode - * whether or not this is a leaf node, whether or not we should default to the - * more distant child in case of missing values, and whether or not this node - * is categorical. The remaining bits are used to encode the feature index for - * this node. Thus, uint8_t may be used for 2**(8 - 3) = 32 or fewer features, - * uint16_t for 2**(16 - 3) = 8192 or fewer, and uint32_t for 536870912 or - * fewer features. - * - * @tparam offset_t An integral type used to indicate the offset from - * this node to its most distant child. This type must be large enough to store - * the largest such offset in the forest model. - */ -template -struct alignas(detail::get_node_alignment()) - node { - // @brief An alias for layout_v - auto constexpr static const layout = layout_v; - // @brief An alias for threshold_t - using threshold_type = threshold_t; - // @brief An alias for index_t - using index_type = index_t; - /* @brief A union to hold either a threshold value or index - * - * All nodes will need EITHER a threshold value, an output value, OR an index - * to data elsewhere that wil be used either for evaluating the node (e.g. an - * index to a categorical set) or creating an output (e.g. an index to vector - * leaf output). This union allows us to store either of these values without - * using additional space for the unused value. - */ - union value_type { - threshold_t value; - index_t index; - }; - /// @brief An alias for metadata_storage_t - using metadata_storage_type = metadata_storage_t; - /// @brief An alias for offset_t - using offset_type = offset_t; - - // TODO(wphicks): Add custom type to ensure given child offset is at least - // one - - // Assumption: Node construction occurs on the host. This allows us to perform - // bound check on the 'feature' parameter. - constexpr node(threshold_type value = threshold_type{}, - bool is_leaf_node = true, - bool default_to_distant_child = false, - bool is_categorical_node = false, - metadata_storage_type feature = metadata_storage_type{}, - offset_type distant_child_offset = offset_type{}) - : aligned_data{ - .inner_data = { - {.value = value}, - distant_child_offset, - construct_metadata(is_leaf_node, default_to_distant_child, is_categorical_node, feature)}} - { - } - - constexpr node(index_type index, - bool is_leaf_node = true, - bool default_to_distant_child = false, - bool is_categorical_node = false, - metadata_storage_type feature = metadata_storage_type{}, - offset_type distant_child_offset = offset_type{}) - : aligned_data{ - .inner_data = { - {.index = index}, - distant_child_offset, - construct_metadata(is_leaf_node, default_to_distant_child, is_categorical_node, feature)}} - { - } - - /* The index of the feature for this node */ - HOST DEVICE auto constexpr feature_index() const - { - return aligned_data.inner_data.metadata & FEATURE_MASK; - } - /* Whether or not this node is a leaf node */ - HOST DEVICE auto constexpr is_leaf() const - { - return !bool(aligned_data.inner_data.distant_offset); - } - /* Whether or not to default to distant child in case of missing values */ - HOST DEVICE auto constexpr default_distant() const - { - return bool(aligned_data.inner_data.metadata & DEFAULT_DISTANT_MASK); - } - /* Whether or not this node is a categorical node */ - HOST DEVICE auto constexpr is_categorical() const - { - return bool(aligned_data.inner_data.metadata & CATEGORICAL_MASK); - } - /* The offset to the child of this node if it evaluates to given condition */ - HOST DEVICE auto constexpr child_offset(bool condition) const - { - if constexpr (layout == tree_layout::depth_first) { - return offset_type{1} + condition * (aligned_data.inner_data.distant_offset - offset_type{1}); - } else if constexpr (layout == tree_layout::breadth_first || - layout == tree_layout::layered_children_together) { - return condition * offset_type{1} + (aligned_data.inner_data.distant_offset - offset_type{1}); - } else { - static_assert(layout == tree_layout::depth_first); - } - } - /* The threshold value for this node */ - HOST DEVICE auto constexpr threshold() const - { - return aligned_data.inner_data.stored_value.value; - } - - /* The index value for this node */ - HOST DEVICE auto const& index() const { return aligned_data.inner_data.stored_value.index; } - /* The output value for this node - * - * @tparam output_t The expected output type for this node. - */ - template - HOST DEVICE auto constexpr output() const - { - if constexpr (has_vector_leaves) { - return aligned_data.inner_data.stored_value.index; - } else { - return aligned_data.inner_data.stored_value.value; - } - } - - private: - /* Define all bit masks required to extract information from the stored - * metadata. The first bit tells us whether or not this is a leaf node, the - * second tells us whether or not we should default to the distant child in - * the case of a missing value, and the third tells us whether or not this is - * a categorical node. The remaining bits indicate the index of the feature - * for this node */ - auto constexpr static const LEAF_BIT = - metadata_storage_type(index_type(sizeof(metadata_storage_type) * 8 - 1)); - auto constexpr static const LEAF_MASK = metadata_storage_type(1 << LEAF_BIT); - auto constexpr static const DEFAULT_DISTANT_BIT = metadata_storage_type(LEAF_BIT - 1); - auto constexpr static const DEFAULT_DISTANT_MASK = - metadata_storage_type(1 << DEFAULT_DISTANT_BIT); - auto constexpr static const CATEGORICAL_BIT = metadata_storage_type(DEFAULT_DISTANT_BIT - 1); - auto constexpr static const CATEGORICAL_MASK = metadata_storage_type(1 << CATEGORICAL_BIT); - auto constexpr static const FEATURE_MASK = - metadata_storage_type(~(LEAF_MASK | DEFAULT_DISTANT_MASK | CATEGORICAL_MASK)); - - // Helper function for bit packing with the above masks - auto static constexpr construct_metadata(bool is_leaf_node = true, - bool default_to_distant_child = false, - bool is_categorical_node = false, - metadata_storage_type feature = metadata_storage_type{}) - { - // Ensure that 'feature' is not truncated. - static_assert(std::is_unsigned_v, "metadata storage must be unsigned"); - if (feature > FEATURE_MASK) { - throw model_import_error{std::string{"The 'feature' value in the node must be at most "} + - std::to_string(FEATURE_MASK) + "."}; - } - - return metadata_storage_type( - (is_leaf_node << LEAF_BIT) + (default_to_distant_child << DEFAULT_DISTANT_BIT) + - (is_categorical_node << CATEGORICAL_BIT) + (feature & FEATURE_MASK)); - } - - auto static constexpr const byte_size = - detail::get_node_alignment(); - - struct inner_data_type { - value_type stored_value; - // TODO (wphicks): It may be possible to store both of the following together - // to save bytes - offset_type distant_offset; - metadata_storage_type metadata; - }; - union aligned_data_type { - inner_data_type inner_data; - char spacer_data[byte_size]; - }; - - aligned_data_type aligned_data; -}; - -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/postprocessor.hpp b/cpp/include/cuml/fil/detail/postprocessor.hpp deleted file mode 100644 index 48cbb370c2..0000000000 --- a/cpp/include/cuml/fil/detail/postprocessor.hpp +++ /dev/null @@ -1,236 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include -#include -#include -#include - -#include - -#include -#include - -#ifndef __CUDACC__ -#include -#endif - -namespace ML { -namespace fil { - -/* Convert the postprocessing operations into a single value - * representing what must be done in the inference kernel - */ -HOST DEVICE inline auto constexpr ops_to_val(row_op row_wise, element_op elem_wise) -{ - return (static_cast>(row_wise) | - static_cast>(elem_wise)); -} - -/* - * Perform postprocessing on raw forest output - * - * @param val Pointer to the raw forest output - * @param output_count The number of output values per row - * @param bias Pointer to bias vector, which is added to the output - * as part of the postprocessing step. The bias vector should have - * the same length as output_count. - * @param out Pointer to the output buffer - * @param stride Number of elements between the first element that must be - * summed for a particular output element and the next. This is typically - * equal to the number of "groves" of trees over which the computation - * was divided. - * @param average_factor The factor by which to divide during the - * normalization step of postprocessing - * @param constant If the postprocessing operation requires a constant, - * it can be passed here. - */ -template -HOST DEVICE void postprocess(infer_kind infer_type, - io_t* val, - index_type output_count, - const io_t* bias, - io_t* out, - index_type stride = index_type{1}, - io_t average_factor = io_t{1}, - io_t constant = io_t{1}) -{ - const bool use_bias = infer_type == infer_kind::default_kind; -#pragma GCC diagnostic push -#pragma GCC diagnostic ignored "-Wunused-but-set-variable" - auto max_index = index_type{}; - auto max_value = std::numeric_limits::lowest(); -#pragma GCC diagnostic pop - for (auto output_index = index_type{}; output_index < output_count; ++output_index) { - auto workspace_index = output_index * stride; - // Add the bias term if use_bias is true. - // The following expression is written to avoid branching. - val[workspace_index] = - val[workspace_index] / average_factor + - bias[output_index * static_cast(use_bias)] * static_cast(use_bias); - if constexpr (elem_wise_v == element_op::signed_square) { - val[workspace_index] = - copysign(val[workspace_index] * val[workspace_index], val[workspace_index]); - } else if constexpr (elem_wise_v == element_op::hinge) { - val[workspace_index] = io_t(val[workspace_index] > io_t{}); - } else if constexpr (elem_wise_v == element_op::sigmoid) { - val[workspace_index] = io_t{1} / (io_t{1} + exp(-constant * val[workspace_index])); - } else if constexpr (elem_wise_v == element_op::exponential) { - val[workspace_index] = exp(val[workspace_index] / constant); - } else if constexpr (elem_wise_v == element_op::logarithm_one_plus_exp) { - val[workspace_index] = log1p(exp(val[workspace_index] / constant)); - } - if constexpr (row_wise_v == row_op::softmax || row_wise_v == row_op::max_index) { - auto is_new_max = val[workspace_index] > max_value; - max_index = is_new_max * output_index + (!is_new_max) * max_index; - max_value = is_new_max * val[workspace_index] + (!is_new_max) * max_value; - } - } - - if constexpr (row_wise_v == row_op::max_index) { - *out = max_index; - } else { -#pragma GCC diagnostic push -#pragma GCC diagnostic ignored "-Wunused-but-set-variable" - auto softmax_normalization = io_t{}; -#pragma GCC diagnostic pop - if constexpr (row_wise_v == row_op::softmax) { - for (auto workspace_index = index_type{}; workspace_index < output_count * stride; - workspace_index += stride) { - val[workspace_index] = exp(val[workspace_index] - max_value); - softmax_normalization += val[workspace_index]; - } - } - - for (auto output_index = index_type{}; output_index < output_count; ++output_index) { - auto workspace_index = output_index * stride; - if constexpr (row_wise_v == row_op::softmax) { - out[output_index] = val[workspace_index] / softmax_normalization; - } else { - out[output_index] = val[workspace_index]; - } - } - } -} - -/* - * Struct which holds all data necessary to perform postprocessing on raw - * output of a forest model - * - * @tparam io_t The type used for input and output to/from the model - * (typically float/double) - * @param row_wise Enum value representing the row-wise post-processing - * operation to perform on the output - * @param elem_wise Enum value representing the element-wise post-processing - * operation to perform on the output - * @param average_factor The factor by which to divide during the - * normalization step of postprocessing - * @param constant If the postprocessing operation requires a constant, - * it can be passed here. - */ -template -struct postprocessor { - HOST DEVICE postprocessor(row_op row_wise = row_op::disable, - element_op elem_wise = element_op::disable, - io_t average_factor = io_t{1}, - io_t constant = io_t{1}) - : average_factor_{average_factor}, - constant_{constant}, - row_wise_{row_wise}, - elem_wise_{elem_wise} - { - } - - HOST DEVICE void operator()(infer_kind infer_type, - io_t* val, - index_type output_count, - const io_t* bias, - io_t* out, - index_type stride = index_type{1}) const - { - switch (ops_to_val(row_wise_, elem_wise_)) { - case ops_to_val(row_op::disable, element_op::signed_square): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::disable, element_op::hinge): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::disable, element_op::sigmoid): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::disable, element_op::exponential): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::disable, element_op::logarithm_one_plus_exp): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::softmax, element_op::disable): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::softmax, element_op::signed_square): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::softmax, element_op::hinge): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::softmax, element_op::sigmoid): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::softmax, element_op::exponential): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::softmax, element_op::logarithm_one_plus_exp): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::max_index, element_op::disable): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::max_index, element_op::signed_square): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::max_index, element_op::hinge): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::max_index, element_op::sigmoid): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::max_index, element_op::exponential): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - case ops_to_val(row_op::max_index, element_op::logarithm_one_plus_exp): - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - break; - default: - postprocess( - infer_type, val, output_count, bias, out, stride, average_factor_, constant_); - } - } - - private: - io_t average_factor_; - io_t constant_; - row_op row_wise_; - element_op elem_wise_; -}; -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/raft_proto/buffer.hpp b/cpp/include/cuml/fil/detail/raft_proto/buffer.hpp deleted file mode 100644 index e7db75c1fe..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/buffer.hpp +++ /dev/null @@ -1,390 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include -#include -#include -#include -#include - -namespace raft_proto { -/** - * @brief A container which may or may not own its own data on host or device - * - */ -template -struct buffer { - using index_type = std::size_t; - using value_type = T; - - using data_store = std::variant, - non_owning_buffer, - owning_buffer, - owning_buffer>; - - buffer() : device_{}, data_{}, size_{}, cached_ptr{nullptr} {} - - /** Construct non-initialized owning buffer */ - buffer(index_type size, - device_type mem_type = device_type::cpu, - int device = 0, - cuda_stream stream = 0) - : device_{[mem_type, &device]() { - auto result = device_id_variant{}; - switch (mem_type) { - case device_type::cpu: result = device_id{device}; break; - case device_type::gpu: result = device_id{device}; break; - } - return result; - }()}, - data_{[this, mem_type, size, stream]() { - auto result = data_store{}; - switch (mem_type) { - case device_type::cpu: result = owning_buffer{size}; break; - case device_type::gpu: - result = owning_buffer{std::get<1>(device_), size, stream}; - break; - } - return result; - }()}, - size_{size}, - cached_ptr{[this]() { - auto result = static_cast(nullptr); - switch (data_.index()) { - case 0: result = std::get<0>(data_).get(); break; - case 1: result = std::get<1>(data_).get(); break; - case 2: result = std::get<2>(data_).get(); break; - case 3: result = std::get<3>(data_).get(); break; - } - return result; - }()} - { - } - - /** Construct non-owning buffer */ - buffer(T* input_data, index_type size, device_type mem_type = device_type::cpu, int device = 0) - : device_{[mem_type, &device]() { - auto result = device_id_variant{}; - switch (mem_type) { - case device_type::cpu: result = device_id{device}; break; - case device_type::gpu: result = device_id{device}; break; - } - return result; - }()}, - data_{[input_data, mem_type]() { - auto result = data_store{}; - switch (mem_type) { - case device_type::cpu: result = non_owning_buffer{input_data}; break; - case device_type::gpu: result = non_owning_buffer{input_data}; break; - } - return result; - }()}, - size_{size}, - cached_ptr{[this]() { - auto result = static_cast(nullptr); - switch (data_.index()) { - case 0: result = std::get<0>(data_).get(); break; - case 1: result = std::get<1>(data_).get(); break; - case 2: result = std::get<2>(data_).get(); break; - case 3: result = std::get<3>(data_).get(); break; - } - return result; - }()} - { - } - - /** - * @brief Construct one buffer from another in the given memory location - * (either on host or on device) - * A buffer constructed in this way is owning and will copy the data from - * the original location - */ - buffer(buffer const& other, - device_type mem_type, - int device = 0, - cuda_stream stream = cuda_stream{}) - : device_{[mem_type, &device]() { - auto result = device_id_variant{}; - switch (mem_type) { - case device_type::cpu: result = device_id{device}; break; - case device_type::gpu: result = device_id{device}; break; - } - return result; - }()}, - data_{[this, &other, mem_type, stream]() { - auto result = data_store{}; - auto result_data = static_cast(nullptr); - if (mem_type == device_type::cpu) { - auto buf = owning_buffer(other.size()); - result_data = buf.get(); - result = std::move(buf); - } else if (mem_type == device_type::gpu) { - auto buf = owning_buffer(std::get<1>(device_), other.size(), stream); - result_data = buf.get(); - result = std::move(buf); - } - copy(result_data, other.data(), other.size(), mem_type, other.memory_type(), stream); - return result; - }()}, - size_{other.size()}, - cached_ptr{[this]() { - auto result = static_cast(nullptr); - switch (data_.index()) { - case 0: result = std::get<0>(data_).get(); break; - case 1: result = std::get<1>(data_).get(); break; - case 2: result = std::get<2>(data_).get(); break; - case 3: result = std::get<3>(data_).get(); break; - } - return result; - }()} - { - } - - /** - * @brief Create owning copy of existing buffer with given stream - * The memory type of this new buffer will be the same as the original - */ - buffer(buffer const& other, cuda_stream stream = cuda_stream{}) - : buffer(other, other.memory_type(), other.device_index(), stream) - { - } - - /** - * @brief Create owning copy of existing buffer - * The memory type of this new buffer will be the same as the original - */ - friend void swap(buffer& first, buffer& second) - { - using std::swap; - swap(first.device_, second.device_); - swap(first.data_, second.data_); - swap(first.size_, second.size_); - swap(first.cached_ptr, second.cached_ptr); - } - buffer& operator=(buffer const& other) - { - auto copy = other; - swap(*this, copy); - return *this; - } - - /** - * @brief Move from existing buffer unless a copy is necessary based on - * memory location - */ - buffer(buffer&& other, device_type mem_type, int device, cuda_stream stream) - : device_{[mem_type, &device]() { - auto result = device_id_variant{}; - switch (mem_type) { - case device_type::cpu: result = device_id{device}; break; - case device_type::gpu: result = device_id{device}; break; - } - return result; - }()}, - data_{[&other, mem_type, device, stream]() { - auto result = data_store{}; - if (mem_type == other.memory_type() && device == other.device_index()) { - result = std::move(other.data_); - } else { - auto* result_data = static_cast(nullptr); - if (mem_type == device_type::cpu) { - auto buf = owning_buffer{other.size()}; - result_data = buf.get(); - result = std::move(buf); - } else if (mem_type == device_type::gpu) { - auto buf = owning_buffer{device, other.size(), stream}; - result_data = buf.get(); - result = std::move(buf); - } - copy(result_data, other.data(), other.size(), mem_type, other.memory_type(), stream); - } - return result; - }()}, - size_{other.size()}, - cached_ptr{[this]() { - auto result = static_cast(nullptr); - switch (data_.index()) { - case 0: result = std::get<0>(data_).get(); break; - case 1: result = std::get<1>(data_).get(); break; - case 2: result = std::get<2>(data_).get(); break; - case 3: result = std::get<3>(data_).get(); break; - } - return result; - }()} - { - } - buffer(buffer&& other, device_type mem_type, int device) - : buffer{std::move(other), mem_type, device, cuda_stream{}} - { - } - buffer(buffer&& other, device_type mem_type) - : buffer{std::move(other), mem_type, 0, cuda_stream{}} - { - } - - buffer(buffer&& other) noexcept - : buffer{std::move(other), other.memory_type(), other.device_index(), cuda_stream{}} - { - } - buffer& operator=(buffer&& other) noexcept - { - data_ = std::move(other.data_); - device_ = std::move(other.device_); - size_ = std::move(other.size_); - cached_ptr = std::move(other.cached_ptr); - return *this; - } - - template < - typename iter_t, - typename = decltype(*std::declval(), void(), ++std::declval(), void())> - buffer(iter_t const& begin, iter_t const& end) - : buffer{static_cast(std::distance(begin, end))} - { - auto index = std::size_t{}; - std::for_each(begin, end, [&index, this](auto&& val) { data()[index++] = val; }); - } - - template < - typename iter_t, - typename = decltype(*std::declval(), void(), ++std::declval(), void())> - buffer(iter_t const& begin, iter_t const& end, device_type mem_type) - : buffer{buffer{begin, end}, mem_type} - { - } - - template < - typename iter_t, - typename = decltype(*std::declval(), void(), ++std::declval(), void())> - buffer(iter_t const& begin, - iter_t const& end, - device_type mem_type, - int device, - cuda_stream stream = cuda_stream{}) - : buffer{buffer{begin, end}, mem_type, device, stream} - { - } - - auto size() const noexcept { return size_; } - HOST DEVICE auto* data() const noexcept { return cached_ptr; } - auto memory_type() const noexcept - { - auto result = device_type{}; - if (device_.index() == 0) { - result = device_type::cpu; - } else { - result = device_type::gpu; - } - return result; - } - - auto device() const noexcept { return device_; } - - auto device_index() const noexcept - { - auto result = int{}; - switch (device_.index()) { - case 0: result = std::get<0>(device_).value(); break; - case 1: result = std::get<1>(device_).value(); break; - } - return result; - } - ~buffer() = default; - - private: - device_id_variant device_; - data_store data_; - index_type size_; - T* cached_ptr; -}; - -template -const_agnostic_same_t copy(buffer& dst, - buffer const& src, - typename buffer::index_type dst_offset, - typename buffer::index_type src_offset, - typename buffer::index_type size, - cuda_stream stream) -{ - if constexpr (bounds_check) { - if (src_offset > src.size() || src.size() - src_offset < size || dst_offset > dst.size() || - dst.size() - dst_offset < size) { - throw out_of_bounds("Attempted copy to or from buffer of inadequate size"); - } - } - copy(dst.data() + dst_offset, - src.data() + src_offset, - size, - dst.memory_type(), - src.memory_type(), - stream); -} - -template -const_agnostic_same_t copy(buffer& dst, buffer const& src, cuda_stream stream) -{ - copy(dst, src, 0, 0, src.size(), stream); -} -template -const_agnostic_same_t copy(buffer& dst, buffer const& src) -{ - copy(dst, src, 0, 0, src.size(), cuda_stream{}); -} - -template -const_agnostic_same_t copy(buffer&& dst, - buffer&& src, - typename buffer::index_type dst_offset, - typename buffer::index_type src_offset, - typename buffer::index_type size, - cuda_stream stream) -{ - if constexpr (bounds_check) { - if (src_offset > src.size() || src.size() - src_offset < size || dst_offset > dst.size() || - dst.size() - dst_offset < size) { - throw out_of_bounds("Attempted copy to or from buffer of inadequate size"); - } - } - copy(dst.data() + dst_offset, - src.data() + src_offset, - size, - dst.memory_type(), - src.memory_type(), - stream); -} - -template -const_agnostic_same_t copy(buffer&& dst, - buffer&& src, - typename buffer::index_type dst_offset, - cuda_stream stream) -{ - copy(dst, src, dst_offset, 0, src.size(), stream); -} - -template -const_agnostic_same_t copy(buffer&& dst, buffer&& src, cuda_stream stream) -{ - copy(dst, src, 0, 0, src.size(), stream); -} -template -const_agnostic_same_t copy(buffer&& dst, buffer&& src) -{ - copy(dst, src, 0, 0, src.size(), cuda_stream{}); -} - -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/ceildiv.hpp b/cpp/include/cuml/fil/detail/raft_proto/ceildiv.hpp deleted file mode 100644 index 429b8af64a..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/ceildiv.hpp +++ /dev/null @@ -1,17 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -#include - -namespace raft_proto { -template -HOST DEVICE auto constexpr ceildiv(T dividend, U divisor) -{ - static_assert(std::is_integral_v && std::is_integral_v, "Arguments must be integers"); - return dividend / divisor + (dividend % divisor != 0); -} -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/cuda_check.hpp b/cpp/include/cuml/fil/detail/raft_proto/cuda_check.hpp deleted file mode 100644 index 8d19d2168a..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/cuda_check.hpp +++ /dev/null @@ -1,19 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#ifdef CUML_ENABLE_GPU -#include -#endif -#include -#include - -namespace raft_proto { -template -void cuda_check(error_t const& err) noexcept(!GPU_ENABLED) -{ - detail::cuda_check(err); -} -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/cuda_stream.hpp b/cpp/include/cuml/fil/detail/raft_proto/cuda_stream.hpp deleted file mode 100644 index 56e4fb519f..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/cuda_stream.hpp +++ /dev/null @@ -1,22 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#ifdef CUML_ENABLE_GPU -#include -#endif - -namespace raft_proto { -#ifdef CUML_ENABLE_GPU -using cuda_stream = cudaStream_t; -#else -using cuda_stream = int; -#endif -inline void synchronize(cuda_stream stream) -{ -#ifdef CUML_ENABLE_GPU - cudaStreamSynchronize(stream); -#endif -} -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/const_agnostic.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/const_agnostic.hpp deleted file mode 100644 index f682066822..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/const_agnostic.hpp +++ /dev/null @@ -1,16 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -namespace raft_proto { -template -using const_agnostic_same_t = - std::enable_if_t, std::remove_const_t>, V>; - -template -inline constexpr auto const_agnostic_same_v = - std::is_same_v, std::remove_const_t>; -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/copy.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/copy.hpp deleted file mode 100644 index 1207367563..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/copy.hpp +++ /dev/null @@ -1,83 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include - -#include -#ifdef CUML_ENABLE_GPU -#include -#endif -#include - -namespace raft_proto { -template -void copy(T* dst, T const* src, uint32_t size, uint32_t dst_offset, uint32_t src_offset) -{ - detail::copy(dst + dst_offset, src + src_offset, size, cuda_stream{}); -} - -template -void copy( - T* dst, T const* src, uint32_t size, uint32_t dst_offset, uint32_t src_offset, cuda_stream stream) -{ - detail::copy(dst + dst_offset, src + src_offset, size, stream); -} - -template -void copy(T* dst, T const* src, uint32_t size) -{ - detail::copy(dst, src, size, cuda_stream{}); -} - -template -void copy(T* dst, T const* src, uint32_t size, cuda_stream stream) -{ - detail::copy(dst, src, size, stream); -} - -template -void copy(T* dst, - T const* src, - uint32_t size, - device_type dst_type, - device_type src_type, - uint32_t dst_offset, - uint32_t src_offset, - cuda_stream stream) -{ - if (dst_type == device_type::gpu && src_type == device_type::gpu) { - detail::copy( - dst + dst_offset, src + src_offset, size, stream); - } else if (dst_type == device_type::cpu && src_type == device_type::cpu) { - detail::copy( - dst + dst_offset, src + src_offset, size, stream); - } else if (dst_type == device_type::gpu && src_type == device_type::cpu) { - detail::copy( - dst + dst_offset, src + src_offset, size, stream); - } else if (dst_type == device_type::cpu && src_type == device_type::gpu) { - detail::copy( - dst + dst_offset, src + src_offset, size, stream); - } -} - -template -void copy(T* dst, T const* src, uint32_t size, device_type dst_type, device_type src_type) -{ - copy(dst, src, size, dst_type, src_type, 0, 0, cuda_stream{}); -} - -template -void copy(T* dst, - T const* src, - uint32_t size, - device_type dst_type, - device_type src_type, - cuda_stream stream) -{ - copy(dst, src, size, dst_type, src_type, 0, 0, stream); -} - -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/copy/cpu.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/copy/cpu.hpp deleted file mode 100644 index db09b985d7..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/copy/cpu.hpp +++ /dev/null @@ -1,39 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include - -#include - -#include -#include - -namespace raft_proto { -namespace detail { - -template -std::enable_if_t, - std::bool_constant>, - void> -copy(T* dst, T const* src, uint32_t size, cuda_stream stream) -{ - std::copy(src, src + size, dst); -} - -template -std::enable_if_t< - std::conjunction_v, - std::bool_constant>, - std::bool_constant>, - void> -copy(T* dst, T const* src, uint32_t size, cuda_stream stream) -{ - throw gpu_unsupported("Copying from or to device in non-GPU build"); -} - -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/copy/gpu.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/copy/gpu.hpp deleted file mode 100644 index 7d4aa15523..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/copy/gpu.hpp +++ /dev/null @@ -1,31 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include - -#include - -#include - -#include - -namespace raft_proto { -namespace detail { - -template -std::enable_if_t< - std::conjunction_v, - std::bool_constant>, - std::bool_constant>, - void> -copy(T* dst, T const* src, uint32_t size, cuda_stream stream) -{ - raft_proto::cuda_check(cudaMemcpyAsync(dst, src, size * sizeof(T), cudaMemcpyDefault, stream)); -} - -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/cuda_check/base.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/cuda_check/base.hpp deleted file mode 100644 index 9a4e928e9b..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/cuda_check/base.hpp +++ /dev/null @@ -1,17 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -namespace raft_proto { -namespace detail { - -template -void cuda_check(error_t const& err) -{ -} - -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/cuda_check/gpu.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/cuda_check/gpu.hpp deleted file mode 100644 index 54681ee68d..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/cuda_check/gpu.hpp +++ /dev/null @@ -1,24 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include - -#include -namespace raft_proto { -namespace detail { - -template <> -inline void cuda_check(cudaError_t const& err) noexcept(false) -{ - if (err != cudaSuccess) { - cudaGetLastError(); - throw bad_cuda_call(cudaGetErrorString(err)); - } -} - -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/device_id/base.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/device_id/base.hpp deleted file mode 100644 index da7ea8d0f6..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/device_id/base.hpp +++ /dev/null @@ -1,18 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -namespace raft_proto { -namespace detail { -template -struct device_id { - using value_type = int; - - device_id(value_type device_index) {} - auto value() const { return value_type{}; } -}; -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/device_id/cpu.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/device_id/cpu.hpp deleted file mode 100644 index 6eda93ad6d..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/device_id/cpu.hpp +++ /dev/null @@ -1,23 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include - -namespace raft_proto { -namespace detail { -template <> -struct device_id { - using value_type = int; - device_id() : id_{value_type{}} {}; - device_id(value_type dev_id) : id_{dev_id} {}; - - auto value() const noexcept { return id_; } - - private: - value_type id_; -}; -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/device_id/gpu.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/device_id/gpu.hpp deleted file mode 100644 index f89c33802e..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/device_id/gpu.hpp +++ /dev/null @@ -1,31 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include - -#include - -namespace raft_proto { -namespace detail { -template <> -struct device_id { - using value_type = typename rmm::cuda_device_id::value_type; - device_id() noexcept(false) - : id_{[]() { - auto raw_id = value_type{}; - raft_proto::cuda_check(cudaGetDevice(&raw_id)); - return raw_id; - }()} {}; - device_id(value_type dev_id) noexcept : id_{dev_id} {}; - - auto value() const noexcept { return id_.value(); } - - private: - rmm::cuda_device_id id_; -}; -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/device_setter/base.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/device_setter/base.hpp deleted file mode 100644 index 939cec5472..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/device_setter/base.hpp +++ /dev/null @@ -1,19 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include - -namespace raft_proto { -namespace detail { - -/** Struct for setting current device within a code block */ -template -struct device_setter { - device_setter(device_id device) {} -}; - -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/device_setter/gpu.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/device_setter/gpu.hpp deleted file mode 100644 index d33af78409..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/device_setter/gpu.hpp +++ /dev/null @@ -1,38 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include - -#include - -#include - -namespace raft_proto { -namespace detail { - -/** Struct for setting current device within a code block */ -template <> -struct device_setter { - device_setter(raft_proto::device_id device) noexcept(false) - : prev_device_{[]() { - auto result = int{}; - raft_proto::cuda_check(cudaGetDevice(&result)); - return result; - }()} - { - raft_proto::cuda_check(cudaSetDevice(device.value())); - } - - ~device_setter() { RAFT_CUDA_TRY_NO_THROW(cudaSetDevice(prev_device_.value())); } - - private: - device_id prev_device_; -}; - -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw.hpp deleted file mode 100644 index 532bed1936..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw.hpp +++ /dev/null @@ -1,13 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include - -namespace raft_proto { -template -using host_only_throw = detail::host_only_throw; -} diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw/base.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw/base.hpp deleted file mode 100644 index 5a16943c8e..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw/base.hpp +++ /dev/null @@ -1,19 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -namespace raft_proto { -namespace detail { -template -struct host_only_throw { - template - host_only_throw(Args&&... args) - { - static_assert(host); // Do not allow constexpr branch to compile if !host - } -}; -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw/cpu.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw/cpu.hpp deleted file mode 100644 index 2bbb8b88fc..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/host_only_throw/cpu.hpp +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include - -namespace raft_proto { -namespace detail { -template -struct host_only_throw { - template - host_only_throw(Args&&... args) noexcept(false) - { - throw T{std::forward(args)...}; - } -}; -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/non_owning_buffer.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/non_owning_buffer.hpp deleted file mode 100644 index 5bcb835bab..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/non_owning_buffer.hpp +++ /dev/null @@ -1,12 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include - -namespace raft_proto { -template -using non_owning_buffer = detail::non_owning_buffer; -} diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/non_owning_buffer/base.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/non_owning_buffer/base.hpp deleted file mode 100644 index e67f1c9eff..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/non_owning_buffer/base.hpp +++ /dev/null @@ -1,28 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -#include -#include - -namespace raft_proto { -namespace detail { -template -struct non_owning_buffer { - // TODO(wphicks): Assess need for buffers of const T - using value_type = std::remove_const_t; - non_owning_buffer() : data_{nullptr} {} - - non_owning_buffer(T* ptr) : data_{ptr} {} - - auto* get() const { return data_; } - - private: - // TODO(wphicks): Back this with RMM-allocated host memory - T* data_; -}; -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer.hpp deleted file mode 100644 index 71aab9d627..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer.hpp +++ /dev/null @@ -1,14 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#ifdef CUML_ENABLE_GPU -#include -#endif -namespace raft_proto { -template -using owning_buffer = detail::owning_buffer; -} diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/base.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/base.hpp deleted file mode 100644 index 553af09ebf..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/base.hpp +++ /dev/null @@ -1,23 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include - -#include - -namespace raft_proto { -namespace detail { - -template -struct owning_buffer { - owning_buffer() {} - owning_buffer(device_id device_id, std::size_t size, cuda_stream stream) {} - auto* get() const { return static_cast(nullptr); } -}; - -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/cpu.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/cpu.hpp deleted file mode 100644 index 130fa9b209..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/cpu.hpp +++ /dev/null @@ -1,31 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include - -#include -#include - -namespace raft_proto { -namespace detail { -template -struct owning_buffer { - // TODO(wphicks): Assess need for buffers of const T - using value_type = std::remove_const_t; - - owning_buffer() : data_{std::unique_ptr{nullptr}} {} - - owning_buffer(std::size_t size) : data_{std::make_unique(size)} {} - - auto* get() const { return data_.get(); } - - private: - // TODO(wphicks): Back this with RMM-allocated host memory - std::unique_ptr data_; -}; -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/gpu.hpp b/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/gpu.hpp deleted file mode 100644 index f3e96f6bdf..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/detail/owning_buffer/gpu.hpp +++ /dev/null @@ -1,41 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include - -#include - -#include - -#include - -namespace raft_proto { -namespace detail { -template -struct owning_buffer { - // TODO(wphicks): Assess need for buffers of const T - using value_type = std::remove_const_t; - owning_buffer() : data_{} {} - - owning_buffer(device_id device_id, - std::size_t size, - cudaStream_t stream) noexcept(false) - : data_{[&device_id, &size, &stream]() { - auto device_context = device_setter{device_id}; - return rmm::device_buffer{size * sizeof(value_type), rmm::cuda_stream_view{stream}}; - }()} - { - } - - auto* get() const { return reinterpret_cast(data_.data()); } - - private: - mutable rmm::device_buffer data_; -}; -} // namespace detail -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/device_id.hpp b/cpp/include/cuml/fil/detail/raft_proto/device_id.hpp deleted file mode 100644 index 8b8fc8165d..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/device_id.hpp +++ /dev/null @@ -1,21 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include -#include -#ifdef CUML_ENABLE_GPU -#include -#endif -#include - -#include - -namespace raft_proto { -template -using device_id = detail::device_id; - -using device_id_variant = std::variant, device_id>; -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/device_setter.hpp b/cpp/include/cuml/fil/detail/raft_proto/device_setter.hpp deleted file mode 100644 index de130ab842..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/device_setter.hpp +++ /dev/null @@ -1,16 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#ifdef CUML_ENABLE_GPU -#include -#endif -#include - -namespace raft_proto { - -using device_setter = detail::device_setter; - -} diff --git a/cpp/include/cuml/fil/detail/raft_proto/device_type.hpp b/cpp/include/cuml/fil/detail/raft_proto/device_type.hpp deleted file mode 100644 index 7adf14aaf0..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/device_type.hpp +++ /dev/null @@ -1,8 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -namespace raft_proto { -enum class device_type { cpu, gpu }; -} diff --git a/cpp/include/cuml/fil/detail/raft_proto/exceptions.hpp b/cpp/include/cuml/fil/detail/raft_proto/exceptions.hpp deleted file mode 100644 index 27329eb8a0..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/exceptions.hpp +++ /dev/null @@ -1,56 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -namespace raft_proto { -struct bad_cuda_call : std::exception { - bad_cuda_call() : bad_cuda_call("CUDA API call failed") {} - bad_cuda_call(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_; } - - private: - char const* msg_; -}; - -struct out_of_bounds : std::exception { - out_of_bounds() : out_of_bounds("Attempted out-of-bounds memory access") {} - out_of_bounds(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_; } - - private: - char const* msg_; -}; - -struct wrong_device_type : std::exception { - wrong_device_type() : wrong_device_type("Attempted to use host data on GPU or device data on CPU") - { - } - wrong_device_type(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_; } - - private: - char const* msg_; -}; - -struct mem_type_mismatch : std::exception { - mem_type_mismatch() : mem_type_mismatch("Memory type does not match expected type") {} - mem_type_mismatch(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_; } - - private: - char const* msg_; -}; - -struct wrong_device : std::exception { - wrong_device() : wrong_device("Attempted to use incorrect device") {} - wrong_device(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_; } - - private: - char const* msg_; -}; - -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/gpu_support.hpp b/cpp/include/cuml/fil/detail/raft_proto/gpu_support.hpp deleted file mode 100644 index b3badd559c..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/gpu_support.hpp +++ /dev/null @@ -1,45 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -#include -#include - -namespace raft_proto { -#ifdef CUML_ENABLE_GPU -auto constexpr static const GPU_ENABLED = true; -#else -auto constexpr static const GPU_ENABLED = false; -#endif - -#ifdef __CUDACC__ -#define HOST __host__ -#define DEVICE __device__ -auto constexpr static const GPU_COMPILATION = true; -#else -#define HOST -#define DEVICE -auto constexpr static const GPU_COMPILATION = false; -#endif - -#ifndef DEBUG -auto constexpr static const DEBUG_ENABLED = false; -#elif DEBUG == 0 -auto constexpr static const DEBUG_ENABLED = false; -#else -auto constexpr static const DEBUG_ENABLED = true; -#endif - -struct gpu_unsupported : std::exception { - gpu_unsupported() : gpu_unsupported("GPU functionality invoked in non-GPU build") {} - gpu_unsupported(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_; } - - private: - char const* msg_; -}; - -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/handle.hpp b/cpp/include/cuml/fil/detail/raft_proto/handle.hpp deleted file mode 100644 index be8e940e62..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/handle.hpp +++ /dev/null @@ -1,43 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -#include -#include -#ifdef CUML_ENABLE_GPU -#include -#endif - -namespace raft_proto { -#ifdef CUML_ENABLE_GPU -struct handle_t { - handle_t(raft::handle_t const* handle_ptr = nullptr) : raft_handle_{handle_ptr} {} - handle_t(raft::handle_t const& raft_handle) : raft_handle_{&raft_handle} {} - auto get_next_usable_stream() const - { - return raft_proto::cuda_stream{raft_handle_->get_next_usable_stream().value()}; - } - auto get_stream_pool_size() const { return raft_handle_->get_stream_pool_size(); } - auto get_usable_stream_count() const { return std::max(get_stream_pool_size(), std::size_t{1}); } - void synchronize() const - { - raft_handle_->sync_stream_pool(); - raft_handle_->sync_stream(); - } - - private: - // Have to store a pointer because handle is not movable - raft::handle_t const* raft_handle_; -}; -#else -struct handle_t { - auto get_next_usable_stream() const { return raft_proto::cuda_stream{}; } - auto get_stream_pool_size() const { return std::size_t{}; } - auto get_usable_stream_count() const { return std::max(get_stream_pool_size(), std::size_t{1}); } - void synchronize() const {} -}; -#endif -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/raft_proto/padding.hpp b/cpp/include/cuml/fil/detail/raft_proto/padding.hpp deleted file mode 100644 index f9f3584b7b..0000000000 --- a/cpp/include/cuml/fil/detail/raft_proto/padding.hpp +++ /dev/null @@ -1,48 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include - -namespace raft_proto { - -/* Return the value that must be added to val to equal the next multiple of - * alignment greater than or equal to val */ -template -HOST DEVICE auto padding_size(T val, U alignment) -{ - auto result = val; - if (alignment != 0) { - auto remainder = val % alignment; - result = alignment - remainder; - result *= (remainder != 0); - } - return result; -} - -/* Return the next multiple of alignment >= val */ -template -HOST DEVICE auto padded_size(T val, U alignment) -{ - return val + padding_size(val, alignment); -} - -/* Return the value that must be added to val to equal the next multiple of - * alignment less than or equal to val */ -template -HOST DEVICE auto downpadding_size(T val, U alignment) -{ - auto result = val; - if (alignment != 0) { result = val % alignment; } - return result; -} - -/* Return the next multiple of alignment <= val */ -template -HOST DEVICE auto downpadded_size(T val, U alignment) -{ - return val - downpadding_size(val, alignment); -} - -} // namespace raft_proto diff --git a/cpp/include/cuml/fil/detail/specialization_types.hpp b/cpp/include/cuml/fil/detail/specialization_types.hpp deleted file mode 100644 index 84bc61c89a..0000000000 --- a/cpp/include/cuml/fil/detail/specialization_types.hpp +++ /dev/null @@ -1,76 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include - -#include -#include -#include -#include - -namespace ML { -namespace fil { -namespace detail { - -/* - * A template used solely to help manage the types which will be compiled in - * standard cuML FIL - * - * The relatively simple and human-readable template parameters of this - * template are translated into the specific types and values required - * to instantiate more complex templates and compile-time checks. - * - * @tparam layout_v The layout of trees within a model - * @tparam double_precision Whether this model should use double-precision - * for floating-point evaluation and 64-bit integers for indexes - * @tparam large_trees Whether this forest expects more than 2**(16 -3) - 1 = - * 8191 features or contains nodes whose child is offset more than 2**16 - 1 = 65535 nodes away. - */ -template -struct specialization_types { - /* The node threshold type to be used based on the template parameters - */ - using threshold_type = std::conditional_t; - /* The type required for specifying indexes to vector leaf outputs or - * non-local categorical data. - */ - using index_type = std::conditional_t; - /* The type used to provide metadata storage for nodes */ - using metadata_type = std::conditional_t; - /* The type used to provide metadata storage for nodes */ - using offset_type = std::conditional_t; - /* The tree layout (alias for layout_v)*/ - auto static constexpr const layout = layout_v; - /* Whether or not this tree requires double precision (alias for - * double_precision) - */ - auto static constexpr const is_double_precision = double_precision; - /* Whether or not this forest contains large trees (alias for - * large_trees) - */ - auto static constexpr const has_large_trees = large_trees; -}; - -/* A variant holding information on all specialization types compiled - * in standard cuML FIL - */ -using specialization_variant = - std::variant, - specialization_types, - specialization_types, - specialization_types, - specialization_types, - specialization_types, - specialization_types, - specialization_types, - specialization_types, - specialization_types, - specialization_types, - specialization_types>; - -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/detail/specializations/device_initialization_macros.hpp b/cpp/include/cuml/fil/detail/specializations/device_initialization_macros.hpp deleted file mode 100644 index c9b29ae4e0..0000000000 --- a/cpp/include/cuml/fil/detail/specializations/device_initialization_macros.hpp +++ /dev/null @@ -1,14 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -/* Declare device initialization function for the types specified by the given - * variant index */ -#define CUML_FIL_INITIALIZE_DEVICE(template_type, variant_index) \ - template_type void \ - initialize_device( \ - raft_proto::device_id); diff --git a/cpp/include/cuml/fil/detail/specializations/forest_macros.hpp b/cpp/include/cuml/fil/detail/specializations/forest_macros.hpp deleted file mode 100644 index 0fa77d12dd..0000000000 --- a/cpp/include/cuml/fil/detail/specializations/forest_macros.hpp +++ /dev/null @@ -1,27 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include - -#include - -/* Macro which, given a variant index, will extract the type of the - * corresponding variant from the specialization_variant type. This allows us - * to specify all forest variants we wish to support in one location and then - * reference them by index elsewhere. */ -#define CUML_FIL_SPEC(variant_index) \ - std::variant_alternative_t - -/* Macro which expands to a full declaration of a forest type corresponding to - * the given variant index. */ -#define CUML_FIL_FOREST(variant_index) \ - forest diff --git a/cpp/include/cuml/fil/detail/specializations/infer_macros.hpp b/cpp/include/cuml/fil/detail/specializations/infer_macros.hpp deleted file mode 100644 index 442c320a8c..0000000000 --- a/cpp/include/cuml/fil/detail/specializations/infer_macros.hpp +++ /dev/null @@ -1,143 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include - -/* Macro which expands to the valid arguments to an inference call for a forest - * model without vector leaves or non-local categorical data.*/ -#define CUML_FIL_SCALAR_LOCAL_ARGS(dev, variant_index) \ - (CUML_FIL_FOREST(variant_index) const&, \ - postprocessor const&, \ - CUML_FIL_SPEC(variant_index)::threshold_type*, \ - CUML_FIL_SPEC(variant_index)::threshold_type*, \ - index_type, \ - index_type, \ - index_type, \ - std::nullptr_t, \ - std::nullptr_t, \ - infer_kind, \ - std::optional, \ - raft_proto::device_id, \ - raft_proto::cuda_stream stream) - -/* Macro which expands to the valid arguments to an inference call for a forest - * model with vector leaves but without non-local categorical data.*/ -#define CUML_FIL_VECTOR_LOCAL_ARGS(dev, variant_index) \ - (CUML_FIL_FOREST(variant_index) const&, \ - postprocessor const&, \ - CUML_FIL_SPEC(variant_index)::threshold_type*, \ - CUML_FIL_SPEC(variant_index)::threshold_type*, \ - index_type, \ - index_type, \ - index_type, \ - CUML_FIL_SPEC(variant_index)::threshold_type*, \ - std::nullptr_t, \ - infer_kind, \ - std::optional, \ - raft_proto::device_id, \ - raft_proto::cuda_stream stream) - -/* Macro which expands to the valid arguments to an inference call for a forest - * model without vector leaves but with non-local categorical data.*/ -#define CUML_FIL_SCALAR_NONLOCAL_ARGS(dev, variant_index) \ - (CUML_FIL_FOREST(variant_index) const&, \ - postprocessor const&, \ - CUML_FIL_SPEC(variant_index)::threshold_type*, \ - CUML_FIL_SPEC(variant_index)::threshold_type*, \ - index_type, \ - index_type, \ - index_type, \ - std::nullptr_t, \ - CUML_FIL_SPEC(variant_index)::index_type*, \ - infer_kind, \ - std::optional, \ - raft_proto::device_id, \ - raft_proto::cuda_stream stream) - -/* Macro which expands to the valid arguments to an inference call for a forest - * model with vector leaves and with non-local categorical data.*/ -#define CUML_FIL_VECTOR_NONLOCAL_ARGS(dev, variant_index) \ - (CUML_FIL_FOREST(variant_index) const&, \ - postprocessor const&, \ - CUML_FIL_SPEC(variant_index)::threshold_type*, \ - CUML_FIL_SPEC(variant_index)::threshold_type*, \ - index_type, \ - index_type, \ - index_type, \ - CUML_FIL_SPEC(variant_index)::threshold_type*, \ - CUML_FIL_SPEC(variant_index)::index_type*, \ - infer_kind, \ - std::optional, \ - raft_proto::device_id, \ - raft_proto::cuda_stream stream) - -/* Macro which expands to the declaration of an inference template for a forest - * of the type indicated by the variant index */ -#define CUML_FIL_INFER_TEMPLATE(template_type, dev, variant_index, categorical) \ - template_type void infer - -/* Macro which expands to the declaration of an inference template for a forest - * of the type indicated by the variant index on the given device type without - * vector leaves or categorical nodes*/ -#define CUML_FIL_INFER_DEV_SCALAR_LEAF_NO_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_TEMPLATE(template_type, dev, variant_index, false) \ - CUML_FIL_SCALAR_LOCAL_ARGS(dev, variant_index); - -/* Macro which expands to the declaration of an inference template for a forest - * of the type indicated by the variant index on the given device type without - * vector leaves and with only local categorical nodes*/ -#define CUML_FIL_INFER_DEV_SCALAR_LEAF_LOCAL_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_TEMPLATE(template_type, dev, variant_index, true) \ - CUML_FIL_SCALAR_LOCAL_ARGS(dev, variant_index); - -/* Macro which expands to the declaration of an inference template for a forest - * of the type indicated by the variant index on the given device type without - * vector leaves and with non-local categorical nodes*/ -#define CUML_FIL_INFER_DEV_SCALAR_LEAF_NONLOCAL_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_TEMPLATE(template_type, dev, variant_index, true) \ - CUML_FIL_SCALAR_NONLOCAL_ARGS(dev, variant_index); - -/* Macro which expands to the declaration of an inference template for a forest - * of the type indicated by the variant index on the given device type with - * vector leaves and without categorical nodes*/ -#define CUML_FIL_INFER_DEV_VECTOR_LEAF_NO_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_TEMPLATE(template_type, dev, variant_index, false) \ - CUML_FIL_VECTOR_LOCAL_ARGS(dev, variant_index); - -/* Macro which expands to the declaration of an inference template for a forest - * of the type indicated by the variant index on the given device type with - * vector leaves and with only local categorical nodes*/ -#define CUML_FIL_INFER_DEV_VECTOR_LEAF_LOCAL_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_TEMPLATE(template_type, dev, variant_index, true) \ - CUML_FIL_VECTOR_LOCAL_ARGS(dev, variant_index); - -/* Macro which expands to the declaration of an inference template for a forest - * of the type indicated by the variant index on the given device type with - * vector leaves and with non-local categorical nodes*/ -#define CUML_FIL_INFER_DEV_VECTOR_LEAF_NONLOCAL_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_TEMPLATE(template_type, dev, variant_index, true) \ - CUML_FIL_VECTOR_NONLOCAL_ARGS(dev, variant_index); - -/* Macro which expands to the declaration of all valid inference templates for - * the given device on the forest type specified by the given variant index */ -#define CUML_FIL_INFER_ALL(template_type, dev, variant_index) \ - CUML_FIL_INFER_DEV_SCALAR_LEAF_NO_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_DEV_SCALAR_LEAF_LOCAL_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_DEV_SCALAR_LEAF_NONLOCAL_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_DEV_VECTOR_LEAF_NO_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_DEV_VECTOR_LEAF_LOCAL_CAT(template_type, dev, variant_index) \ - CUML_FIL_INFER_DEV_VECTOR_LEAF_NONLOCAL_CAT(template_type, dev, variant_index) diff --git a/cpp/include/cuml/fil/exceptions.hpp b/cpp/include/cuml/fil/exceptions.hpp deleted file mode 100644 index 39c00a43e3..0000000000 --- a/cpp/include/cuml/fil/exceptions.hpp +++ /dev/null @@ -1,64 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include - -namespace ML { -namespace fil { - -/** Exception indicating model is incompatible with FIL */ -struct unusable_model_exception : std::exception { - unusable_model_exception() : msg_{"Model is not compatible with FIL"} {} - unusable_model_exception(std::string msg) : msg_{msg} {} - unusable_model_exception(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_.c_str(); } - - private: - std::string msg_; -}; - -/** Exception indicating model import failed */ -struct model_import_error : std::exception { - model_import_error() : model_import_error("Error while importing model") {} - model_import_error(std::string msg) : msg_{msg} {} - model_import_error(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_.c_str(); } - - private: - std::string msg_; -}; - -/** - * Exception indicating a mismatch between the type of input data and the - * model - * - * This typically occurs when doubles are provided as input to a model with - * float thresholds or vice versa. - */ -struct type_error : std::exception { - type_error() : type_error("Model cannot be used with given data type") {} - type_error(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_; } - - private: - char const* msg_; -}; - -/** - * Exception indicating a runtime error. - */ -struct runtime_error : std::exception { - runtime_error() : runtime_error("Runtime error") {} - runtime_error(char const* msg) : msg_{msg} {} - runtime_error(std::string const& msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_.c_str(); } - - private: - std::string msg_; -}; - -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/forest_model.hpp b/cpp/include/cuml/fil/forest_model.hpp deleted file mode 100644 index e91a6de265..0000000000 --- a/cpp/include/cuml/fil/forest_model.hpp +++ /dev/null @@ -1,306 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include -#include -#include -#include - -#include - -#include -#include -#include - -namespace ML { -namespace fil { - -/** - * A model used for performing inference with FIL - * - * This struct is a wrapper for all variants of decision_forest supported by a - * standard FIL build. - */ -struct forest_model { - /** Wrap a decision_forest in a full forest_model object */ - forest_model(decision_forest_variant&& forest = decision_forest_variant{}) - : decision_forest_{forest} - { - } - - /** The number of features per row expected by the model */ - auto num_features() - { - return std::visit([](auto&& concrete_forest) { return concrete_forest.num_features(); }, - decision_forest_); - } - - /** The number of outputs per row generated by the model */ - auto num_outputs() - { - return std::visit([](auto&& concrete_forest) { return concrete_forest.num_outputs(); }, - decision_forest_); - } - - /** The number of trees in the model */ - auto num_trees() - { - return std::visit([](auto&& concrete_forest) { return concrete_forest.num_trees(); }, - decision_forest_); - } - - /** Whether or not leaf nodes use vector outputs */ - auto has_vector_leaves() - { - return std::visit([](auto&& concrete_forest) { return concrete_forest.has_vector_leaves(); }, - decision_forest_); - } - - /** The operation used for postprocessing all outputs for a single row */ - auto row_postprocessing() - { - return std::visit([](auto&& concrete_forest) { return concrete_forest.row_postprocessing(); }, - decision_forest_); - } - - /** Setter for row_postprocessing() */ - void set_row_postprocessing(row_op val) - { - return std::visit( - [&val](auto&& concrete_forest) { concrete_forest.set_row_postprocessing(val); }, - decision_forest_); - } - - /** The operation used for postprocessing each element of the output for a - * single row */ - auto elem_postprocessing() - { - return std::visit([](auto&& concrete_forest) { return concrete_forest.elem_postprocessing(); }, - decision_forest_); - } - - /** The type of memory (device/host) where the model is stored */ - auto memory_type() - { - return std::visit([](auto&& concrete_forest) { return concrete_forest.memory_type(); }, - decision_forest_); - } - - /** The ID of the device on which this model is loaded */ - auto device_index() - { - return std::visit([](auto&& concrete_forest) { return concrete_forest.device_index(); }, - decision_forest_); - } - - /** Whether or not model is loaded at double precision */ - auto is_double_precision() - { - return std::visit( - [](auto&& concrete_forest) { - return std::is_same_v::io_type, - double>; - }, - decision_forest_); - } - - /** - * Perform inference on given input - * - * @param[out] output The buffer where model output should be stored. - * This must be of size at least ROWS x num_outputs(). - * @param[in] input The buffer containing input data. - * @param[in] stream A raft_proto::cuda_stream, which (on GPU-enabled builds) is - * a transparent wrapper for the cudaStream_t or (on CPU-only builds) a - * CUDA-free placeholder object. - * @param[in] predict_type Type of inference to perform. Defaults to summing - * the outputs of all trees and produce an output per row. If set to - * "per_tree", we will instead output all outputs of individual trees. - * If set to "leaf_id", we will output the integer ID of the leaf node - * for each tree. - * @param[in] specified_chunk_size: Specifies the mini-batch size for - * processing. This has different meanings on CPU and GPU, but on GPU it - * corresponds to the number of rows evaluated per inference iteration - * on a single block. It can take on any power of 2 from 1 to 32, and - * runtime performance is quite sensitive to the value chosen. In general, - * larger batches benefit from higher values, but it is hard to predict the - * optimal value a priori. If omitted, a heuristic will be used to select a - * reasonable value. On CPU, this argument can generally just be omitted. - */ - template - void predict(raft_proto::buffer& output, - raft_proto::buffer const& input, - raft_proto::cuda_stream stream = raft_proto::cuda_stream{}, - infer_kind predict_type = infer_kind::default_kind, - std::optional specified_chunk_size = std::nullopt) - { - std::visit( - [this, predict_type, &output, &input, &stream, &specified_chunk_size]( - auto&& concrete_forest) { - if constexpr (std::is_same_v< - typename std::remove_reference_t::io_type, - io_t>) { - concrete_forest.predict(output, input, stream, predict_type, specified_chunk_size); - } else { - throw type_error("Input type does not match model_type"); - } - }, - decision_forest_); - } - - /** - * Perform inference on given input - * - * @param[in] handle The raft_proto::handle_t (wrapper for raft::handle_t - * on GPU) which will be used to provide streams for evaluation. - * @param[out] output The buffer where model output should be stored. If - * this buffer is on host while the model is on device or vice versa, - * work will be distributed across available streams to copy the data back - * to this output location. This must be of size at least ROWS x num_outputs(). - * @param[in] input The buffer containing input data. If - * this buffer is on host while the model is on device or vice versa, - * work will be distributed across available streams to copy the input data - * to the appropriate location and perform inference. - * @param[in] predict_type Type of inference to perform. Defaults to summing - * the outputs of all trees and produce an output per row. If set to - * "per_tree", we will instead output all outputs of individual trees. - * If set to "leaf_id", we will output the integer ID of the leaf node - * for each tree. - * @param[in] specified_chunk_size: Specifies the mini-batch size for - * processing. This has different meanings on CPU and GPU, but on GPU it - * corresponds to the number of rows evaluated per inference iteration - * on a single block. It can take on any power of 2 from 1 to 32, and - * runtime performance is quite sensitive to the value chosen. In general, - * larger batches benefit from higher values, but it is hard to predict the - * optimal value a priori. If omitted, a heuristic will be used to select a - * reasonable value. On CPU, this argument can generally just be omitted. - */ - template - void predict(raft_proto::handle_t const& handle, - raft_proto::buffer& output, - raft_proto::buffer const& input, - infer_kind predict_type = infer_kind::default_kind, - std::optional specified_chunk_size = std::nullopt) - { - std::visit( - [this, predict_type, &handle, &output, &input, &specified_chunk_size]( - auto&& concrete_forest) { - using model_io_t = typename std::remove_reference_t::io_type; - if constexpr (std::is_same_v) { - if (output.memory_type() == memory_type() && input.memory_type() == memory_type()) { - concrete_forest.predict( - output, input, handle.get_next_usable_stream(), predict_type, specified_chunk_size); - } else { - auto constexpr static const MIN_CHUNKS_PER_PARTITION = std::size_t{64}; - auto constexpr static const MAX_CHUNK_SIZE = std::size_t{64}; - - auto row_count = input.size() / num_features(); - auto partition_size = - std::max(raft_proto::ceildiv(row_count, handle.get_usable_stream_count()), - specified_chunk_size.value_or(MAX_CHUNK_SIZE) * MIN_CHUNKS_PER_PARTITION); - auto partition_count = raft_proto::ceildiv(row_count, partition_size); - for (auto i = std::size_t{}; i < partition_count; ++i) { - auto stream = handle.get_next_usable_stream(); - auto rows_in_this_partition = - std::min(partition_size, row_count - i * partition_size); - auto partition_in = raft_proto::buffer{}; - if (input.memory_type() != memory_type()) { - partition_in = - raft_proto::buffer{rows_in_this_partition * num_features(), memory_type()}; - raft_proto::copy(partition_in, - input, - 0, - i * partition_size * num_features(), - partition_in.size(), - stream); - } else { - partition_in = - raft_proto::buffer{input.data() + i * partition_size * num_features(), - rows_in_this_partition * num_features(), - memory_type()}; - } - auto partition_out = raft_proto::buffer{}; - if (output.memory_type() != memory_type()) { - partition_out = - raft_proto::buffer{rows_in_this_partition * num_outputs(), memory_type()}; - } else { - partition_out = - raft_proto::buffer{output.data() + i * partition_size * num_outputs(), - rows_in_this_partition * num_outputs(), - memory_type()}; - } - concrete_forest.predict( - partition_out, partition_in, stream, predict_type, specified_chunk_size); - if (output.memory_type() != memory_type()) { - raft_proto::copy(output, - partition_out, - i * partition_size * num_outputs(), - 0, - partition_out.size(), - stream); - } - } - } - } else { - throw type_error("Input type does not match model_type"); - } - }, - decision_forest_); - } - - /** - * Perform inference on given input - * - * @param[in] handle The raft_proto::handle_t (wrapper for raft::handle_t - * on GPU) which will be used to provide streams for evaluation. - * @param[out] output Pointer to the memory location where output should end - * up - * @param[in] input Pointer to the input data - * @param[in] num_rows Number of rows in input - * @param[in] out_mem_type The memory type (device/host) of the output - * buffer - * @param[in] in_mem_type The memory type (device/host) of the input buffer - * @param[in] predict_type Type of inference to perform. Defaults to summing - * the outputs of all trees and produce an output per row. If set to - * "per_tree", we will instead output all outputs of individual trees. - * If set to "leaf_id", we will output the integer ID of the leaf node - * for each tree. - * @param[in] specified_chunk_size: Specifies the mini-batch size for - * processing. This has different meanings on CPU and GPU, but on GPU it - * corresponds to the number of rows evaluated per inference iteration - * on a single block. It can take on any power of 2 from 1 to 32, and - * runtime performance is quite sensitive to the value chosen. In general, - * larger batches benefit from higher values, but it is hard to predict the - * optimal value a priori. If omitted, a heuristic will be used to select a - * reasonable value. On CPU, this argument can generally just be omitted. - */ - template - void predict(raft_proto::handle_t const& handle, - io_t* output, - io_t* input, - std::size_t num_rows, - raft_proto::device_type out_mem_type, - raft_proto::device_type in_mem_type, - infer_kind predict_type = infer_kind::default_kind, - std::optional specified_chunk_size = std::nullopt) - { - int current_device_id; - raft_proto::cuda_check(cudaGetDevice(¤t_device_id)); - auto out_buffer = - raft_proto::buffer{output, num_rows * num_outputs(), out_mem_type, current_device_id}; - auto in_buffer = - raft_proto::buffer{input, num_rows * num_features(), in_mem_type, current_device_id}; - predict(handle, out_buffer, in_buffer, predict_type, specified_chunk_size); - } - - private: - decision_forest_variant decision_forest_; -}; - -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/infer_kind.hpp b/cpp/include/cuml/fil/infer_kind.hpp deleted file mode 100644 index 386f34775f..0000000000 --- a/cpp/include/cuml/fil/infer_kind.hpp +++ /dev/null @@ -1,10 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -namespace ML { -namespace fil { -enum class infer_kind : unsigned char { default_kind = 0, per_tree = 1, leaf_id = 2 }; -} -} // namespace ML diff --git a/cpp/include/cuml/fil/postproc_ops.hpp b/cpp/include/cuml/fil/postproc_ops.hpp deleted file mode 100644 index 66b4461ef9..0000000000 --- a/cpp/include/cuml/fil/postproc_ops.hpp +++ /dev/null @@ -1,27 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -namespace ML { -namespace fil { - -/** Enum representing possible row-wise operations on output */ -enum struct row_op : unsigned char { - disable = 0b00100000, - softmax = 0b01000000, - max_index = 0b10000000 -}; - -/** Enum representing possible element-wise operations on output */ -enum struct element_op : unsigned char { - disable = 0b00000000, - signed_square = 0b00000001, - hinge = 0b00000010, - sigmoid = 0b00000100, - exponential = 0b00001000, - logarithm_one_plus_exp = 0b00010000 -}; - -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/fil/tree_layout.hpp b/cpp/include/cuml/fil/tree_layout.hpp deleted file mode 100644 index e926e66602..0000000000 --- a/cpp/include/cuml/fil/tree_layout.hpp +++ /dev/null @@ -1,19 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -namespace ML { -namespace fil { -enum class tree_layout : unsigned char { - depth_first = 0, - breadth_first = 1, - // Traverse forest by proceeding through the root nodes of each tree first, - // followed by the hot and distant children of those root nodes for each tree, - // and so forth. This traversal order ensures that all nodes of a - // particular tree at a particular depth are traversed together. - layered_children_together = 2 -}; - -} -} // namespace ML diff --git a/cpp/include/cuml/fil/treelite_importer.hpp b/cpp/include/cuml/fil/treelite_importer.hpp deleted file mode 100644 index 853fc5d39b..0000000000 --- a/cpp/include/cuml/fil/treelite_importer.hpp +++ /dev/null @@ -1,509 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include -#include -#include -#include -#include - -#include -#include - -namespace ML { -namespace fil { - -namespace detail { - -struct postproc_params_t { - element_op element = element_op::disable; - row_op row = row_op::disable; - double constant = 1.0; -}; -} // namespace detail - -/** - * Struct used to import a model from Treelite to FIL - * - * @tparam layout The in-memory layout for nodes to be loaded into FIL - */ -template -struct treelite_importer { - auto static constexpr const traversal_order = []() constexpr { - if constexpr (layout == tree_layout::depth_first) { - return ML::forest::forest_order::depth_first; - } else if constexpr (layout == tree_layout::breadth_first) { - return ML::forest::forest_order::breadth_first; - } else if constexpr (layout == tree_layout::layered_children_together) { - return ML::forest::forest_order::layered_children_together; - } else { - static_assert(layout == tree_layout::depth_first, - "Layout not yet implemented in treelite importer for FIL"); - } - }(); - - auto get_node_count(treelite::Model const& tl_model) - { - return ML::forest::tree_accumulate( - tl_model, index_type{}, [](auto&& count, auto&& tree) { return count + tree.num_nodes; }); - } - - /* Return vector of offsets between each node and its most distant child */ - auto get_offsets(treelite::Model const& tl_model) - { - auto node_count = get_node_count(tl_model); - auto result = std::vector(node_count); - auto parent_indexes = std::vector{}; - parent_indexes.reserve(node_count); - ML::forest::node_transform( - tl_model, - std::back_inserter(parent_indexes), - [](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { return parent_index; }); - for (auto i = std::size_t{}; i < node_count; ++i) { - result[parent_indexes[i]] = i - parent_indexes[i]; - } - return result; - } - - auto num_trees(treelite::Model const& tl_model) - { - auto result = index_type{}; - std::visit([&result](auto&& concrete_tl_model) { result = concrete_tl_model.trees.size(); }, - tl_model.variant_); - return result; - } - - auto get_tree_sizes(treelite::Model const& tl_model) - { - auto result = std::vector{}; - tree_transform( - tl_model, std::back_inserter(result), [](auto&& tree) { return tree.num_nodes; }); - return result; - } - - auto get_num_class(treelite::Model const& tl_model) - { - return static_cast(tl_model.num_class[0]); - } - - auto get_num_feature(treelite::Model const& tl_model) - { - return static_cast(tl_model.num_feature); - } - - auto get_max_num_categories(treelite::Model const& tl_model) - { - return ML::forest::node_accumulate( - tl_model, - index_type{}, - [](auto&& cur_accum, auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - return std::max(cur_accum, static_cast(node.max_num_categories())); - }); - } - - auto get_num_categorical_nodes(treelite::Model const& tl_model) - { - return ML::forest::node_accumulate( - tl_model, - index_type{}, - [](auto&& cur_accum, auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - return cur_accum + static_cast(node.is_categorical()); - }); - } - - auto get_num_leaf_vector_nodes(treelite::Model const& tl_model) - { - return ML::forest::node_accumulate( - tl_model, - index_type{}, - [](auto&& cur_accum, auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - auto accum = cur_accum; - if (node.is_leaf() && node.get_output().size() > 1) { ++accum; } - return accum; - }); - } - - auto get_average_factor(treelite::Model const& tl_model) - { - auto result = double{}; - if (tl_model.average_tree_output) { - if (tl_model.task_type == treelite::TaskType::kMultiClf && - tl_model.leaf_vector_shape[1] == 1) { // grove-per-class - result = num_trees(tl_model) / tl_model.num_class[0]; - } else { - result = num_trees(tl_model); - } - } else { - result = 1.0; - } - return result; - } - - auto get_bias(treelite::Model const& tl_model) { return tl_model.base_scores.AsVector(); } - - auto get_postproc_params(treelite::Model const& tl_model) - { - auto result = detail::postproc_params_t{}; - auto tl_pred_transform = tl_model.postprocessor; - if (tl_pred_transform == std::string{"identity"} || - tl_pred_transform == std::string{"identity_multiclass"}) { - result.element = element_op::disable; - result.row = row_op::disable; - } else if (tl_pred_transform == std::string{"signed_square"}) { - result.element = element_op::signed_square; - } else if (tl_pred_transform == std::string{"hinge"}) { - result.element = element_op::hinge; - } else if (tl_pred_transform == std::string{"sigmoid"}) { - result.constant = tl_model.sigmoid_alpha; - result.element = element_op::sigmoid; - } else if (tl_pred_transform == std::string{"exponential"}) { - result.element = element_op::exponential; - } else if (tl_pred_transform == std::string{"exponential_standard_ratio"}) { - result.constant = -tl_model.ratio_c / std::log(2); - result.element = element_op::exponential; - } else if (tl_pred_transform == std::string{"logarithm_one_plus_exp"}) { - result.element = element_op::logarithm_one_plus_exp; - } else if (tl_pred_transform == std::string{"max_index"}) { - result.row = row_op::max_index; - } else if (tl_pred_transform == std::string{"softmax"}) { - result.row = row_op::softmax; - } else if (tl_pred_transform == std::string{"multiclass_ova"}) { - result.constant = tl_model.sigmoid_alpha; - result.element = element_op::sigmoid; - } else { - throw unusable_model_exception{"Unrecognized Treelite pred_transform string"}; - } - return result; - } - - auto uses_double_thresholds(treelite::Model const& tl_model) - { - auto result = false; - switch (tl_model.GetThresholdType()) { - case treelite::TypeInfo::kFloat64: result = true; break; - case treelite::TypeInfo::kFloat32: result = false; break; - default: throw unusable_model_exception("Unrecognized Treelite threshold type"); - } - return result; - } - - auto uses_double_outputs(treelite::Model const& tl_model) - { - auto result = false; - switch (tl_model.GetThresholdType()) { - case treelite::TypeInfo::kFloat64: result = true; break; - case treelite::TypeInfo::kFloat32: result = false; break; - case treelite::TypeInfo::kUInt32: result = false; break; - default: throw unusable_model_exception("Unrecognized Treelite threshold type"); - } - return result; - } - - auto uses_integer_outputs(treelite::Model const& tl_model) - { - auto result = false; - switch (tl_model.GetThresholdType()) { - case treelite::TypeInfo::kFloat64: result = false; break; - case treelite::TypeInfo::kFloat32: result = false; break; - case treelite::TypeInfo::kUInt32: result = true; break; - default: throw unusable_model_exception("Unrecognized Treelite threshold type"); - } - return result; - } - - /** - * Assuming that the correct decision_forest variant has been - * identified, import to that variant - */ - template - auto import_to_specific_variant(index_type target_variant_index, - treelite::Model const& tl_model, - index_type num_class, - index_type num_feature, - index_type max_num_categories, - std::vector const& offsets, - index_type align_bytes = index_type{}, - raft_proto::device_type mem_type = raft_proto::device_type::cpu, - int device = 0, - raft_proto::cuda_stream stream = raft_proto::cuda_stream{}) - { - auto result = decision_forest_variant{}; - if constexpr (variant_index != std::variant_size_v) { - if (variant_index == target_variant_index) { - using forest_model_t = std::variant_alternative_t; - if constexpr (traversal_order == ML::forest::forest_order::layered_children_together) { - // Cannot align whole trees with layered traversal order, since trees - // are mingled together - align_bytes = index_type{}; - } - auto builder = - detail::decision_forest_builder(max_num_categories, align_bytes); - auto node_index = index_type{}; - ML::forest::node_for_each( - tl_model, - [&builder, &offsets, &node_index]( - auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - try { - if (node.is_leaf()) { - auto output = node.get_output(); - builder.set_output_size(output.size()); - if (output.size() > index_type{1}) { - builder.add_leaf_vector_node( - std::begin(output), std::end(output), node.get_treelite_id(), depth); - } else { - builder.add_node(typename forest_model_t::io_type(output[0]), - node.get_treelite_id(), - depth, - true); - } - } else { - if (node.is_categorical()) { - auto categories = node.get_categories(); - builder.add_categorical_node(std::begin(categories), - std::end(categories), - node.get_treelite_id(), - depth, - node.default_distant(), - node.get_feature(), - offsets[node_index]); - } else { - builder.add_node(typename forest_model_t::threshold_type(node.threshold()), - node.get_treelite_id(), - depth, - false, - node.default_distant(), - false, - node.get_feature(), - offsets[node_index], - node.is_inclusive()); - } - } - } catch (const model_import_error& e) { - throw model_import_error{std::string{"Tree "} + std::to_string(tree_id) + ", Node " + - std::to_string(node.get_treelite_id()) + ": " + e.what()}; - } - ++node_index; - }); - - builder.set_average_factor(get_average_factor(tl_model)); - builder.set_bias(get_bias(tl_model)); - auto postproc_params = get_postproc_params(tl_model); - builder.set_element_postproc(postproc_params.element); - builder.set_row_postproc(postproc_params.row); - builder.set_postproc_constant(postproc_params.constant); - - result.template emplace( - builder.get_decision_forest(num_feature, num_class, mem_type, device, stream)); - } else { - result = import_to_specific_variant(target_variant_index, - tl_model, - num_class, - num_feature, - max_num_categories, - offsets, - align_bytes, - mem_type, - device, - stream); - } - } - return result; - } - - /** - * Import a treelite model to FIL - * - * Load a model from Treelite to a FIL forest_model. The model will be - * inspected to determine the correct underlying decision_forest variant to - * use within the forest_model object. - * - * @param tl_model The Treelite Model to load - * @param align_bytes If non-zero, ensure that each tree is stored in a - * multiple of this value of bytes by padding with empty nodes. This can - * be useful for increasing the likelihood that successive reads will take - * place within a single cache line. On GPU, a value of 128 can be used for - * this purpose. On CPU, a value of either 0 or 64 typically produces - * optimal performance. - * @param use_double_precision Whether or not to use 64 bit floats for model - * evaluation and 64 bit ints for applicable indexing - * @param dev_type Which device type to use for inference (CPU or GPU) - * @param device For GPU execution, the device id for the device on which this - * model is to be loaded - * @param stream The CUDA stream to use for loading this model (can be - * omitted for CPU). - */ - forest_model import(treelite::Model const& tl_model, - index_type align_bytes = index_type{}, - std::optional use_double_precision = std::nullopt, - raft_proto::device_type dev_type = raft_proto::device_type::cpu, - int device = 0, - raft_proto::cuda_stream stream = raft_proto::cuda_stream{}) - { - // Handle degenerate trees (a single root node with no child) - if (auto processed_tl_model = detail::convert_degenerate_trees(tl_model); processed_tl_model) { - return import( - *processed_tl_model.get(), align_bytes, use_double_precision, dev_type, device, stream); - } - - ASSERT(tl_model.num_target == 1, "FIL does not support multi-target model"); - // Check tree annotation (assignment) - if (tl_model.task_type == treelite::TaskType::kMultiClf) { - // Must be either vector leaf or grove-per-class - if (tl_model.leaf_vector_shape[1] > 1) { // vector-leaf - ASSERT(tl_model.leaf_vector_shape[1] == int(tl_model.num_class[0]), - "Vector leaf must be equal to num_class = %d", - tl_model.num_class[0]); - auto tree_count = num_trees(tl_model); - for (decltype(tree_count) tree_id = 0; tree_id < tree_count; ++tree_id) { - ASSERT(tl_model.class_id[tree_id] == -1, "Tree %d has invalid class assignment", tree_id); - } - } else { // grove-per-class - auto tree_count = num_trees(tl_model); - for (decltype(tree_count) tree_id = 0; tree_id < tree_count; ++tree_id) { - ASSERT(tl_model.class_id[tree_id] == int(tree_id % tl_model.num_class[0]), - "Tree %d has invalid class assignment", - tree_id); - } - } - } - - auto result = decision_forest_variant{}; - auto num_feature = get_num_feature(tl_model); - auto max_num_categories = get_max_num_categories(tl_model); - auto num_categorical_nodes = get_num_categorical_nodes(tl_model); - auto num_leaf_vector_nodes = get_num_leaf_vector_nodes(tl_model); - auto use_double_thresholds = use_double_precision.value_or(uses_double_thresholds(tl_model)); - - auto offsets = get_offsets(tl_model); - auto max_offset = *std::max_element(std::begin(offsets), std::end(offsets)); - - auto variant_index = get_forest_variant_index(use_double_thresholds, - max_offset, - num_feature, - num_categorical_nodes, - max_num_categories, - num_leaf_vector_nodes, - layout); - auto num_class = get_num_class(tl_model); - return forest_model{import_to_specific_variant(variant_index, - tl_model, - num_class, - num_feature, - max_num_categories, - offsets, - align_bytes, - dev_type, - device, - stream)}; - } -}; - -/** - * Import a treelite model to FIL - * - * Load a model from Treelite to a FIL forest_model. The model will be - * inspected to determine the correct underlying decision_forest variant to - * use within the forest_model object. - * - * @param tl_model The Treelite Model to load - * @param layout The in-memory layout of nodes in the loaded forest - * @param align_bytes If non-zero, ensure that each tree is stored in a - * multiple of this value of bytes by padding with empty nodes. This can - * be useful for increasing the likelihood that successive reads will take - * place within a single cache line. On GPU, a value of 128 can be used for - * this purpose. On CPU, a value of either 0 or 64 typically produces - * optimal performance. - * @param use_double_precision Whether or not to use 64 bit floats for model - * evaluation and 64 bit ints for applicable indexing - * @param dev_type Which device type to use for inference (CPU or GPU) - * @param device For GPU execution, the device id for the device on which this - * model is to be loaded - * @param stream The CUDA stream to use for loading this model (can be - * omitted for CPU). - */ -inline auto import_from_treelite_model( - treelite::Model const& tl_model, - tree_layout layout = preferred_tree_layout, - index_type align_bytes = index_type{}, - std::optional use_double_precision = std::nullopt, - raft_proto::device_type dev_type = raft_proto::device_type::cpu, - int device = 0, - raft_proto::cuda_stream stream = raft_proto::cuda_stream{}) -{ - auto result = forest_model{}; - switch (layout) { - case tree_layout::depth_first: - result = treelite_importer{}.import( - tl_model, align_bytes, use_double_precision, dev_type, device, stream); - break; - case tree_layout::breadth_first: - result = treelite_importer{}.import( - tl_model, align_bytes, use_double_precision, dev_type, device, stream); - break; - case tree_layout::layered_children_together: - result = treelite_importer{}.import( - tl_model, align_bytes, use_double_precision, dev_type, device, stream); - break; - } - return result; -} - -/** - * Import a treelite model handle to FIL - * - * Load a model from a Treelite model handle (type-erased treelite::Model - * object) to a FIL forest_model. The model will be inspected to determine the - * correct underlying decision_forest variant to use within the forest_model - * object. - * - * @param tl_handle The Treelite ModelHandle to load - * @param layout The in-memory layout of nodes in the loaded forest - * @param align_bytes If non-zero, ensure that each tree is stored in a - * multiple of this value of bytes by padding with empty nodes. This can - * be useful for increasing the likelihood that successive reads will take - * place within a single cache line. On GPU, a value of 128 can be used for - * this purpose. On CPU, a value of either 0 or 64 typically produces - * optimal performance. - * @param use_double_precision Whether or not to use 64 bit floats for model - * evaluation and 64 bit ints for applicable indexing - * @param dev_type Which device type to use for inference (CPU or GPU) - * @param device For GPU execution, the device id for the device on which this - * model is to be loaded - * @param stream The CUDA stream to use for loading this model (can be - * omitted for CPU). - */ -inline auto import_from_treelite_handle( - TreeliteModelHandle tl_handle, - tree_layout layout = preferred_tree_layout, - index_type align_bytes = index_type{}, - std::optional use_double_precision = std::nullopt, - raft_proto::device_type dev_type = raft_proto::device_type::cpu, - int device = 0, - raft_proto::cuda_stream stream = raft_proto::cuda_stream{}) -{ - return import_from_treelite_model(*static_cast(tl_handle), - layout, - align_bytes, - use_double_precision, - dev_type, - device, - stream); -} - -} // namespace fil -} // namespace ML diff --git a/cpp/include/cuml/forest/README.md b/cpp/include/cuml/forest/README.md deleted file mode 100644 index 9955dc1de6..0000000000 --- a/cpp/include/cuml/forest/README.md +++ /dev/null @@ -1,7 +0,0 @@ -# Forest Primitives - -This directory contains headers which are useful for dealing with forest models. - -## Directories - -- `traversal`: This directory contains tools for traversing nodes of a forest in a variety of ways. These tools are not intended for high-speed traversal but for systematically extracting information from existing forest-like data structures in order to facilitate high-speed forest operations. diff --git a/cpp/include/cuml/forest/exceptions.hpp b/cpp/include/cuml/forest/exceptions.hpp deleted file mode 100644 index f876c5bee1..0000000000 --- a/cpp/include/cuml/forest/exceptions.hpp +++ /dev/null @@ -1,21 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include - -namespace ML { -namespace forest { -struct traversal_exception : std::exception { - traversal_exception() : msg_{"Error encountered while traversing forest"} {} - traversal_exception(std::string msg) : msg_{msg} {} - traversal_exception(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_.c_str(); } - - private: - std::string msg_; -}; -} // namespace forest -} // namespace ML diff --git a/cpp/include/cuml/forest/integrations/treelite.hpp b/cpp/include/cuml/forest/integrations/treelite.hpp deleted file mode 100644 index 29984dfe8d..0000000000 --- a/cpp/include/cuml/forest/integrations/treelite.hpp +++ /dev/null @@ -1,226 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include - -#include - -#include -#include -#include - -namespace ML { -namespace forest { - -using TREELITE_NODE_ID_T = int; - -template -struct treelite_traversal_node : public traversal_node { - treelite_traversal_node(treelite::Tree const& tl_tree, - id_type node_id) - : traversal_node{}, tl_tree_{tl_tree}, node_id_{node_id} - { - } - - bool is_leaf() const override { return tl_tree_.IsLeaf(node_id_); } - - id_type hot_child() const override - { - auto result = id_type{}; - if (left_is_hot()) { - result = tl_tree_.LeftChild(node_id_); - } else { - result = tl_tree_.RightChild(node_id_); - } - return result; - } - - id_type distant_child() const override - { - auto result = id_type{}; - if (left_is_hot()) { - result = tl_tree_.RightChild(node_id_); - } else { - result = tl_tree_.LeftChild(node_id_); - } - return result; - } - - auto default_distant() const { return tl_tree_.DefaultChild(node_id_) == distant_child(); } - - auto get_feature() const { return tl_tree_.SplitIndex(node_id_); } - - auto is_inclusive() const - { - auto tl_operator = tl_tree_.ComparisonOp(node_id_); - return tl_operator == treelite::Operator::kGT || tl_operator == treelite::Operator::kLE; - } - - auto is_categorical() const - { - return tl_tree_.NodeType(node_id_) == treelite::TreeNodeType::kCategoricalTestNode; - } - - auto get_categories() const { return tl_tree_.CategoryList(node_id_); } - - auto threshold() const { return tl_tree_.Threshold(node_id_); } - -// Temporarily disable free-nonheap-object warning to work around spurious warnings emitted by -// GCC 14.x. See https://github.com/rapidsai/cuml/pull/7471#issuecomment-3525796585 for more -// details. -// TODO(hcho3): Remove this pragma once GCC is upgraded to 15. -#pragma GCC diagnostic push -#pragma GCC diagnostic ignored "-Wfree-nonheap-object" - auto max_num_categories() const - { - auto result = std::remove_const_t>{}; - if (is_categorical()) { - auto categories = get_categories(); - if (categories.size() != 0) { - result = *std::max_element(std::begin(categories), std::end(categories)) + 1; - } - } - return result; - } -#pragma GCC diagnostic pop - - auto get_output() const - { - auto result = std::vector{}; - if (tl_tree_.HasLeafVector(node_id_)) { - result = tl_tree_.LeafVector(node_id_); - } else { - result.push_back(tl_tree_.LeafValue(node_id_)); - } - return result; - } - - auto get_treelite_id() const { return node_id_; } - - private: - treelite::Tree const& tl_tree_; - id_type node_id_; - - auto left_is_hot() const - { - auto result = false; - if (is_categorical()) { - if (tl_tree_.CategoryListRightChild(node_id_)) { result = true; } - } else { - auto tl_operator = tl_tree_.ComparisonOp(node_id_); - if (tl_operator == treelite::Operator::kLT || tl_operator == treelite::Operator::kLE) { - result = false; - } else if (tl_operator == treelite::Operator::kGT || tl_operator == treelite::Operator::kGE) { - result = true; - } else { - throw traversal_exception("Unrecognized Treelite operator"); - } - } - return result; - } -}; - -template -struct treelite_traversal_forest - : public traversal_forest> { - private: - using base_type = traversal_forest>; - - public: - using node_type = typename base_type::node_type; - using node_id_type = typename base_type::node_id_type; - using tree_id_type = typename base_type::tree_id_type; - using node_uid_type = typename base_type::node_uid_type; - - treelite_traversal_forest(treelite::ModelPreset const& tl_model) - : traversal_forest>{[&tl_model]() { - auto result = std::vector{}; - result.reserve(tl_model.GetNumTree()); - for (auto i = std::size_t{}; i < tl_model.GetNumTree(); ++i) { - result.push_back(std::make_pair(i, TREELITE_NODE_ID_T{})); - } - return result; - }()}, - tl_model_{tl_model} - { - } - - node_type get_node(tree_id_type tree_id, node_id_type node_id) const override - { - return node_type{tl_model_.trees[tree_id], node_id}; - } - - private: - treelite::ModelPreset const& tl_model_; -}; - -template -void tree_for_each(treelite::Model const& tl_model, lambda_t&& lambda) -{ - std::visit( - [&lambda](auto&& concrete_tl_model) { - std::for_each(std::begin(concrete_tl_model.trees), std::end(concrete_tl_model.trees), lambda); - }, - tl_model.variant_); -} - -template -void tree_transform(treelite::Model const& tl_model, iter_t out_iter, lambda_t&& lambda) -{ - std::visit( - [&lambda, out_iter](auto&& concrete_tl_model) { - std::transform( - std::begin(concrete_tl_model.trees), std::end(concrete_tl_model.trees), out_iter, lambda); - }, - tl_model.variant_); -} - -template -auto tree_accumulate(treelite::Model const& tl_model, T init, lambda_t&& lambda) -{ - return std::visit( - [&lambda, init](auto&& concrete_tl_model) { - return std::accumulate( - std::begin(concrete_tl_model.trees), std::end(concrete_tl_model.trees), init, lambda); - }, - tl_model.variant_); -} - -template -void node_for_each(treelite::Model const& tl_model, lambda_t&& lambda) -{ - std::visit( - [&lambda](auto&& concrete_tl_model) { - treelite_traversal_forest{concrete_tl_model}.template for_each(lambda); - }, - tl_model.variant_); -} - -template -void node_transform(treelite::Model const& tl_model, iter_t output_iter, lambda_t&& lambda) -{ - node_for_each( - tl_model, - [&output_iter, &lambda](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - *output_iter = lambda(tree_id, node, depth, parent_index); - ++output_iter; - }); -} - -template -auto node_accumulate(treelite::Model const& tl_model, T init, lambda_t&& lambda) -{ - auto result = init; - node_for_each( - tl_model, [&result, &lambda](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - result = lambda(result, tree_id, node, depth, parent_index); - }); - return result; -} - -} // namespace forest -} // namespace ML diff --git a/cpp/include/cuml/forest/traversal/traversal_forest.hpp b/cpp/include/cuml/forest/traversal/traversal_forest.hpp deleted file mode 100644 index dbe70e2061..0000000000 --- a/cpp/include/cuml/forest/traversal/traversal_forest.hpp +++ /dev/null @@ -1,190 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include - -#include -#include -#include - -namespace ML { -namespace forest { - -namespace detail { -/** A template for storing nodes in order to traverse them in the - * indicated order */ -template -struct traversal_container { - using backing_container_t = - std::conditional_t, std::queue>; - void add(T const& val) { data_.push(val); } - void add(T const& hot, T const& distant) - { - if constexpr (order == forest_order::depth_first) { - data_.push(distant); - data_.push(hot); - } else { - data_.push(hot); - data_.push(distant); - } - } - auto next() - { - if constexpr (std::is_same_v>) { - auto result = data_.top(); - data_.pop(); - return result; - } else { - auto result = data_.front(); - data_.pop(); - return result; - } - } - auto peek() - { - if constexpr (std::is_same_v>) { - return data_.top(); - } else { - return data_.front(); - } - } - [[nodiscard]] auto empty() { return data_.empty(); } - auto size() { return data_.size(); } - - private: - backing_container_t data_; -}; -} // namespace detail - // - -template , typename tree_id_t = std::size_t> -struct traversal_forest { - using node_type = node_t; - using node_id_type = typename node_type::id_type; - using tree_id_type = tree_id_t; - using node_uid_type = std::pair; - using index_type = std::size_t; - - virtual node_type get_node(tree_id_type tree_id, node_id_type node_id) const = 0; - - traversal_forest(std::vector&& root_node_uids) : root_node_uids_{root_node_uids} {} - - template - void for_each(lambda_t&& lambda) const - { - auto to_be_visited = detail::traversal_container< - order, - std::conditional_t>>{}; - auto parent_indices = detail::traversal_container{}; - auto cur_index = index_type{}; - if constexpr (order == forest_order::depth_first || order == forest_order::breadth_first) { - for (auto const& root_node_uid : root_node_uids_) { - to_be_visited.add(std::make_pair(root_node_uid, std::size_t{})); - parent_indices.add(cur_index); - while (!to_be_visited.empty()) { - auto [node_uid, depth] = to_be_visited.next(); - auto parent_index = parent_indices.next(); - auto node = get_node(node_uid); - lambda(node_uid.first, node, depth, parent_index); - if (!node.is_leaf()) { - auto hot_uid = std::make_pair(std::make_pair(node_uid.first, node.hot_child()), - depth + index_type{1}); - auto distant_uid = std::make_pair(std::make_pair(node_uid.first, node.distant_child()), - depth + index_type{1}); - to_be_visited.add(hot_uid, distant_uid); - parent_indices.add(cur_index, cur_index); - } - ++cur_index; - } - } - } else if constexpr (order == forest_order::layered_children_segregated) { - for (auto const& root_node_uid : root_node_uids_) { - to_be_visited.add(root_node_uid); - parent_indices.add(cur_index++); - } - cur_index = index_type{}; - auto depth = index_type{}; - while (!to_be_visited.empty()) { - auto layer_node_uids = std::vector{}; - auto layer_parent_indices = std::vector{}; - while (!to_be_visited.empty()) { - layer_node_uids.push_back(to_be_visited.next()); - layer_parent_indices.push_back(parent_indices.next()); - } - for (auto layer_index = index_type{}; layer_index < layer_node_uids.size(); ++layer_index) { - auto node_uid = layer_node_uids[layer_index]; - auto parent_index = layer_parent_indices[layer_index]; - auto node = get_node(node_uid); - lambda(node_uid.first, node, depth, parent_index); - if (!node.is_leaf()) { - auto hot_uid = std::make_pair(node_uid.first, node.hot_child()); - to_be_visited.add(hot_uid); - parent_indices.add(cur_index); - } - ++cur_index; - } - // Reset cur_index before iterating through distant nodes - cur_index -= layer_node_uids.size(); - for (auto layer_index = index_type{}; layer_index < layer_node_uids.size(); ++layer_index) { - auto node_uid = layer_node_uids[layer_index]; - auto node = get_node(node_uid); - if (!node.is_leaf()) { - auto distant_uid = std::make_pair(node_uid.first, node.distant_child()); - to_be_visited.add(distant_uid); - parent_indices.add(cur_index); - } - ++cur_index; - } - ++depth; - } - } else if constexpr (order == forest_order::layered_children_together) { - for (auto const& root_node_uid : root_node_uids_) { - to_be_visited.add(root_node_uid); - parent_indices.add(cur_index++); - } - cur_index = index_type{}; - auto depth = index_type{}; - while (!to_be_visited.empty()) { - auto layer_node_uids = std::vector{}; - auto layer_parent_indices = std::vector{}; - while (!to_be_visited.empty()) { - layer_node_uids.push_back(to_be_visited.next()); - layer_parent_indices.push_back(parent_indices.next()); - } - for (auto layer_index = index_type{}; layer_index < layer_node_uids.size(); ++layer_index) { - auto node_uid = layer_node_uids[layer_index]; - auto parent_index = layer_parent_indices[layer_index]; - auto node = get_node(node_uid); - lambda(node_uid.first, node, depth, parent_index); - if (!node.is_leaf()) { - auto hot_uid = std::make_pair(node_uid.first, node.hot_child()); - auto distant_uid = std::make_pair(node_uid.first, node.distant_child()); - to_be_visited.add(hot_uid, distant_uid); - parent_indices.add(cur_index, cur_index); - } - ++cur_index; - } - ++depth; - } - } - } - - private: - auto get_node(node_uid_type node_uid) const { return get_node(node_uid.first, node_uid.second); } - - std::vector root_node_uids_{}; -}; - -} // namespace forest -} // namespace ML diff --git a/cpp/include/cuml/forest/traversal/traversal_node.hpp b/cpp/include/cuml/forest/traversal/traversal_node.hpp deleted file mode 100644 index 89280dedba..0000000000 --- a/cpp/include/cuml/forest/traversal/traversal_node.hpp +++ /dev/null @@ -1,39 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once -#include -#include -#include - -namespace ML { -namespace forest { - -/** Exception indicating model is incompatible with FIL */ -struct parentless_node_exception : std::exception { - parentless_node_exception() : msg_{"Node does not track its parent"} {} - parentless_node_exception(std::string msg) : msg_{msg} {} - parentless_node_exception(char const* msg) : msg_{msg} {} - virtual char const* what() const noexcept { return msg_.c_str(); } - - private: - std::string msg_; -}; - -template -struct traversal_node { - public: - using id_type = id_t; - virtual bool is_leaf() const = 0; - virtual id_type hot_child() const = 0; - virtual id_type distant_child() const = 0; - virtual id_type parent() const - { - throw parentless_node_exception(); - return id_type{}; - } -}; - -} // namespace forest -} // namespace ML diff --git a/cpp/include/cuml/forest/traversal/traversal_order.hpp b/cpp/include/cuml/forest/traversal/traversal_order.hpp deleted file mode 100644 index 6c5b2aa85a..0000000000 --- a/cpp/include/cuml/forest/traversal/traversal_order.hpp +++ /dev/null @@ -1,38 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -namespace ML { -namespace forest { - -/* A class used to specify the order in which nodes of a forest should be - * traversed - * - * Because the meaning of "left" and "right" vary by convention, we refer to the two children of a - * node as "hot" or "distant" rather than left or right. The "hot" child is the one which is - * traversed soonest after the parent, and the "distant" child is traversed latest. - */ -enum class forest_order : unsigned char { - // Traverse forest by proceeding depth-first through each tree - // consecutively - depth_first = 0, - // Traverse forest by proceeding breadth-first through each tree - // consecutively - breadth_first = 1, - // Traverse forest by proceeding through the root nodes of each tree first, - // followed by the hot and distant children of those root nodes for each tree, - // and so forth. This traversal order ensures that all nodes of a - // particular tree at a particular depth are traversed together. - layered_children_together = 2, - // Traverse forest by proceeding through the root nodes of each tree first, - // followed by all of the hot children of those root nodes, then all of - // the distant children of those root nodes, and so forth. This - // traversal order ensures that all hot children at a particular depth - // are traversed together, followed by all distant children. - layered_children_segregated = 3 -}; - -} // namespace forest -} // namespace ML diff --git a/cpp/src/fil/infer0.cpp b/cpp/src/fil/infer0.cpp deleted file mode 100644 index de1fdabe3b..0000000000 --- a/cpp/src/fil/infer0.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 0) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer0.cu b/cpp/src/fil/infer0.cu deleted file mode 100644 index da100b1eaf..0000000000 --- a/cpp/src/fil/infer0.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 0) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 0) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer1.cpp b/cpp/src/fil/infer1.cpp deleted file mode 100644 index a8f79b9a4a..0000000000 --- a/cpp/src/fil/infer1.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 1) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer1.cu b/cpp/src/fil/infer1.cu deleted file mode 100644 index 6cad588c9a..0000000000 --- a/cpp/src/fil/infer1.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 1) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 1) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer10.cpp b/cpp/src/fil/infer10.cpp deleted file mode 100644 index 7ebac74ad9..0000000000 --- a/cpp/src/fil/infer10.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 10) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer10.cu b/cpp/src/fil/infer10.cu deleted file mode 100644 index d1f8827111..0000000000 --- a/cpp/src/fil/infer10.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 10) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 10) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer11.cpp b/cpp/src/fil/infer11.cpp deleted file mode 100644 index 8fd11188ef..0000000000 --- a/cpp/src/fil/infer11.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 11) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer11.cu b/cpp/src/fil/infer11.cu deleted file mode 100644 index a641a4d613..0000000000 --- a/cpp/src/fil/infer11.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 11) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 11) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer2.cpp b/cpp/src/fil/infer2.cpp deleted file mode 100644 index c62a2cfc74..0000000000 --- a/cpp/src/fil/infer2.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 2) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer2.cu b/cpp/src/fil/infer2.cu deleted file mode 100644 index bf3e313426..0000000000 --- a/cpp/src/fil/infer2.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 2) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 2) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer3.cpp b/cpp/src/fil/infer3.cpp deleted file mode 100644 index 7e20ad0a0a..0000000000 --- a/cpp/src/fil/infer3.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 3) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer3.cu b/cpp/src/fil/infer3.cu deleted file mode 100644 index 79a33f1692..0000000000 --- a/cpp/src/fil/infer3.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 3) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 3) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer4.cpp b/cpp/src/fil/infer4.cpp deleted file mode 100644 index b0dd9aac95..0000000000 --- a/cpp/src/fil/infer4.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 4) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer4.cu b/cpp/src/fil/infer4.cu deleted file mode 100644 index 28750e6d85..0000000000 --- a/cpp/src/fil/infer4.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 4) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 4) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer5.cpp b/cpp/src/fil/infer5.cpp deleted file mode 100644 index 4efab38c10..0000000000 --- a/cpp/src/fil/infer5.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 5) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer5.cu b/cpp/src/fil/infer5.cu deleted file mode 100644 index 47be1174a7..0000000000 --- a/cpp/src/fil/infer5.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 5) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 5) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer6.cpp b/cpp/src/fil/infer6.cpp deleted file mode 100644 index 9878a78b25..0000000000 --- a/cpp/src/fil/infer6.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 6) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer6.cu b/cpp/src/fil/infer6.cu deleted file mode 100644 index e30aeff8fd..0000000000 --- a/cpp/src/fil/infer6.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 6) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 6) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer7.cpp b/cpp/src/fil/infer7.cpp deleted file mode 100644 index 5ae6e27b40..0000000000 --- a/cpp/src/fil/infer7.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 7) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer7.cu b/cpp/src/fil/infer7.cu deleted file mode 100644 index c6d0ff641c..0000000000 --- a/cpp/src/fil/infer7.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 7) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 7) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer8.cpp b/cpp/src/fil/infer8.cpp deleted file mode 100644 index 9fb92e05ec..0000000000 --- a/cpp/src/fil/infer8.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 8) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer8.cu b/cpp/src/fil/infer8.cu deleted file mode 100644 index 3967981f79..0000000000 --- a/cpp/src/fil/infer8.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 8) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 8) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer9.cpp b/cpp/src/fil/infer9.cpp deleted file mode 100644 index 5390e279d6..0000000000 --- a/cpp/src/fil/infer9.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::cpu, 9) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/src/fil/infer9.cu b/cpp/src/fil/infer9.cu deleted file mode 100644 index 9eac246d83..0000000000 --- a/cpp/src/fil/infer9.cu +++ /dev/null @@ -1,20 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#include -#include -#include -#include -namespace ML { -namespace fil { -namespace detail { -namespace inference { -CUML_FIL_INFER_ALL(template, raft_proto::device_type::gpu, 9) -} -namespace device_initialization { -CUML_FIL_INITIALIZE_DEVICE(template, 9) -} -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 6d918e0e91..563553cf74 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -42,6 +42,7 @@ function(ConfigureTest) CUDA::cusolver${_ctk_static_suffix} CUDA::cusparse${_ctk_static_suffix} $<$:CUDA::cufft${_ctk_static_suffix_cufft}> + $<$:${CUML_NVFOREST_TARGET}> CUDA::cudart_static rmm::rmm raft::raft @@ -104,19 +105,6 @@ if(all_algo OR explainer_algo) ConfigureTest(PREFIX SG NAME SHAP_KERNEL_TEST sg/shap_kernel.cu ML_INCLUDE) endif() -if(all_algo OR fil_algo) - ConfigureTest(PREFIX SG NAME HOST_BUFFER_TEST sg/fil/raft_proto/buffer.cpp ML_INCLUDE) - ConfigureTest(PREFIX SG NAME DEVICE_BUFFER_TEST sg/fil/raft_proto/buffer.cu ML_INCLUDE) - ConfigureTest(PREFIX SG NAME FOREST_TRAVERSAL_TEST sg/forest/traversal_forest.cpp ML_INCLUDE) - ConfigureTest(PREFIX SG NAME TREELITE_TRAVERSAL_TEST sg/forest/treelite_traversal.cpp ML_INCLUDE) - ConfigureTest( - PREFIX SG - NAME TREELITE_IMPORTER_TEST - sg/fil/treelite_importer.cpp sg/fil/treelite_importer_invalid_inputs.cpp - sg/fil/decision_forest_builder_invalid_inputs.cpp ML_INCLUDE - ) -endif() - # todo: organize linear models better if(all_algo OR linearregression_algo diff --git a/cpp/tests/sg/fil/decision_forest_builder_invalid_inputs.cpp b/cpp/tests/sg/fil/decision_forest_builder_invalid_inputs.cpp deleted file mode 100644 index 5e508573a9..0000000000 --- a/cpp/tests/sg/fil/decision_forest_builder_invalid_inputs.cpp +++ /dev/null @@ -1,82 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include -#include -#include -#include - -#include -#include - -#include -#include - -namespace ML { -namespace fil { -namespace detail { - -using test_forest_t = - decision_forest; - -TEST(DecisionForestBuilder, CategoricalStorageOffsetOutOfBounds) -{ - auto builder = decision_forest_builder( - /*max_num_categories=*/std::uint32_t{33}, /*align_bytes=*/std::uint32_t{0}); - - // Construct a malformed categorical node that references non-local category - // storage at an out-of-range offset. This should be rejected by the invariant - // checks in get_decision_forest(). - builder.add_node(std::uint32_t{1234}, - /*tl_node_id=*/0, - /*depth=*/0, - /*is_leaf_node=*/false, - /*default_to_distant_child=*/false, - /*is_categorical_node=*/true, - /*feature=*/0, - /*offset=*/1); - - ASSERT_THAT( - [&] { builder.get_decision_forest(/*num_feature=*/1, /*num_class=*/1); }, - testing::ThrowsMessage(testing::HasSubstr("storage offset out of bounds"))); -} - -TEST(DecisionForestBuilder, CategoricalBitsetExtentOutOfBounds) -{ - auto builder = decision_forest_builder( - /*max_num_categories=*/std::uint32_t{33}, /*align_bytes=*/std::uint32_t{0}); - - // Create a valid categorical node first, which allocates non-local storage: - // categorical_storage_ = [num_categories, packed_bin_data] - std::array categories{0}; - builder.add_categorical_node(categories.begin(), - categories.end(), - /*tl_node_id=*/0, - /*depth=*/0, - /*default_to_distant_child=*/false, - /*feature=*/0, - /*offset=*/1); - - // Construct another categorical node that points at offset=1, i.e. the first - // packed bin entry rather than the metadata entry. The value at offset=1 is - // interpreted as stored_num_cats, and the resulting bins_required exceeds the - // available headroom. - builder.add_node(std::uint32_t{1}, - /*tl_node_id=*/1, - /*depth=*/1, - /*is_leaf_node=*/false, - /*default_to_distant_child=*/false, - /*is_categorical_node=*/true, - /*feature=*/0, - /*offset=*/1); - - ASSERT_THAT([&] { builder.get_decision_forest(/*num_feature=*/1, /*num_class=*/1); }, - testing::ThrowsMessage( - testing::HasSubstr("bitset extends past categorical_storage end"))); -} - -} // namespace detail -} // namespace fil -} // namespace ML diff --git a/cpp/tests/sg/fil/raft_proto/buffer.cpp b/cpp/tests/sg/fil/raft_proto/buffer.cpp deleted file mode 100644 index 137bcccc6f..0000000000 --- a/cpp/tests/sg/fil/raft_proto/buffer.cpp +++ /dev/null @@ -1,379 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include -#include -#include -#include -#include - -#include -#include - -namespace raft_proto { - -TEST(Buffer, default_buffer) -{ - auto buf = buffer(); - EXPECT_EQ(buf.memory_type(), device_type::cpu); - EXPECT_EQ(buf.size(), 0); - EXPECT_EQ(buf.device_index(), 0); -} - -TEST(Buffer, device_buffer) -{ - auto data = std::vector{1, 2, 3}; - auto test_buffers = std::vector>{}; - test_buffers.emplace_back(data.size(), device_type::gpu, 0, cuda_stream{}); - test_buffers.emplace_back(data.size(), device_type::gpu, 0); - test_buffers.emplace_back(data.size(), device_type::gpu); - - for (auto& buf : test_buffers) { - ASSERT_EQ(buf.memory_type(), device_type::gpu); - ASSERT_EQ(buf.size(), data.size()); -#ifdef CUML_ENABLE_GPU - ASSERT_NE(buf.data(), nullptr); - - auto data_out = std::vector(data.size()); - cudaMemcpy(static_cast(buf.data()), - static_cast(data.data()), - sizeof(int) * data.size(), - cudaMemcpyHostToDevice); - cudaMemcpy(static_cast(data_out.data()), - static_cast(buf.data()), - sizeof(int) * data.size(), - cudaMemcpyDeviceToHost); - EXPECT_THAT(data_out, testing::ElementsAreArray(data)); -#endif - } -} - -TEST(Buffer, non_owning_device_buffer) -{ - auto data = std::vector{1, 2, 3}; - auto* ptr_d = static_cast(nullptr); -#ifdef CUML_ENABLE_GPU - cudaMalloc(reinterpret_cast(&ptr_d), sizeof(int) * data.size()); - cudaMemcpy(static_cast(ptr_d), - static_cast(data.data()), - sizeof(int) * data.size(), - cudaMemcpyHostToDevice); -#endif - auto test_buffers = std::vector>{}; - test_buffers.emplace_back(ptr_d, data.size(), device_type::gpu, 0); - test_buffers.emplace_back(ptr_d, data.size(), device_type::gpu); -#ifdef CUML_ENABLE_GPU - - for (auto& buf : test_buffers) { - ASSERT_EQ(buf.memory_type(), device_type::gpu); - ASSERT_EQ(buf.size(), data.size()); - ASSERT_EQ(buf.data(), ptr_d); - - auto data_out = std::vector(data.size()); - cudaMemcpy(static_cast(data_out.data()), - static_cast(buf.data()), - sizeof(int) * data.size(), - cudaMemcpyDeviceToHost); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); - } - cudaFree(reinterpret_cast(ptr_d)); -#endif -} - -TEST(Buffer, host_buffer) -{ - auto data = std::vector{1, 2, 3}; - auto test_buffers = std::vector>{}; - test_buffers.emplace_back(data.size(), device_type::cpu, 0, cuda_stream{}); - test_buffers.emplace_back(data.size(), device_type::cpu, 0); - test_buffers.emplace_back(data.size(), device_type::cpu); - test_buffers.emplace_back(data.size()); - - for (auto& buf : test_buffers) { - ASSERT_EQ(buf.memory_type(), device_type::cpu); - ASSERT_EQ(buf.size(), data.size()); - ASSERT_NE(buf.data(), nullptr); - - std::copy(data.begin(), data.end(), buf.data()); - - auto data_out = std::vector(buf.data(), buf.data() + buf.size()); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); - } -} - -TEST(Buffer, host_buffer_from_iters) -{ - auto data = std::vector{1, 2, 3}; - auto test_buffers = std::vector>{}; - test_buffers.emplace_back(std::begin(data), std::end(data)); - - for (auto& buf : test_buffers) { - ASSERT_EQ(buf.memory_type(), device_type::cpu); - ASSERT_EQ(buf.size(), data.size()); - ASSERT_NE(buf.data(), nullptr); - - std::copy(data.begin(), data.end(), buf.data()); - - auto data_out = std::vector(buf.data(), buf.data() + buf.size()); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); - } -} - -TEST(Buffer, device_buffer_from_iters) -{ - auto data = std::vector{1, 2, 3}; - auto test_buffers = std::vector>{}; - test_buffers.emplace_back(std::begin(data), std::end(data), device_type::gpu); - test_buffers.emplace_back(std::begin(data), std::end(data), device_type::gpu, 0); - test_buffers.emplace_back(std::begin(data), std::end(data), device_type::gpu, 0, cuda_stream{}); - - for (auto& buf : test_buffers) { - ASSERT_EQ(buf.memory_type(), device_type::gpu); - ASSERT_EQ(buf.size(), data.size()); -#ifdef CUML_ENABLE_GPU - ASSERT_NE(buf.data(), nullptr); - - auto data_out = std::vector(data.size()); - cudaMemcpy(static_cast(buf.data()), - static_cast(data.data()), - sizeof(int) * data.size(), - cudaMemcpyHostToDevice); - cudaMemcpy(static_cast(data_out.data()), - static_cast(buf.data()), - sizeof(int) * data.size(), - cudaMemcpyDeviceToHost); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); -#endif - } -} - -TEST(Buffer, non_owning_host_buffer) -{ - auto data = std::vector{1, 2, 3}; - std::vector> test_buffers; - test_buffers.emplace_back(data.data(), data.size(), device_type::cpu, 0); - ASSERT_EQ(test_buffers.back().memory_type(), device_type::cpu); - ASSERT_EQ(test_buffers.back().size(), data.size()); - ASSERT_EQ(test_buffers.back().data(), data.data()); - test_buffers.emplace_back(data.data(), data.size(), device_type::cpu); - ASSERT_EQ(test_buffers.back().memory_type(), device_type::cpu); - ASSERT_EQ(test_buffers.back().size(), data.size()); - ASSERT_EQ(test_buffers.back().data(), data.data()); - test_buffers.emplace_back(data.data(), data.size()); - ASSERT_EQ(test_buffers.back().memory_type(), device_type::cpu); - ASSERT_EQ(test_buffers.back().size(), data.size()); - ASSERT_EQ(test_buffers.back().data(), data.data()); - - for (auto& buf : test_buffers) { - ASSERT_EQ(buf.memory_type(), device_type::cpu); - ASSERT_EQ(buf.size(), data.size()); - ASSERT_EQ(buf.data(), data.data()); - - auto data_out = std::vector(buf.data(), buf.data() + buf.size()); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); - } -} - -TEST(Buffer, copy_buffer) -{ - auto data = std::vector{1, 2, 3}; - auto orig_buffer = buffer(data.data(), data.size(), device_type::cpu); - - auto test_buffers = std::vector>{}; - test_buffers.emplace_back(orig_buffer); - test_buffers.emplace_back(orig_buffer, device_type::cpu); - test_buffers.emplace_back(orig_buffer, device_type::cpu, 0); - test_buffers.emplace_back(orig_buffer, device_type::cpu, 0, cuda_stream{}); - - for (auto& buf : test_buffers) { - ASSERT_EQ(buf.memory_type(), device_type::cpu); - ASSERT_EQ(buf.size(), data.size()); - ASSERT_NE(buf.data(), orig_buffer.data()); - - auto data_out = std::vector(buf.data(), buf.data() + buf.size()); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); - -#ifdef CUML_ENABLE_GPU - auto test_dev_buffers = std::vector>{}; - test_dev_buffers.emplace_back(orig_buffer, device_type::gpu); - test_dev_buffers.emplace_back(orig_buffer, device_type::gpu, 0); - test_dev_buffers.emplace_back(orig_buffer, device_type::gpu, 0, cuda_stream{}); - for (auto& dev_buf : test_dev_buffers) { - data_out = std::vector(data.size()); - cuda_check(cudaMemcpy(static_cast(data_out.data()), - static_cast(dev_buf.data()), - dev_buf.size() * sizeof(int), - cudaMemcpyDefault)); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); - - auto test_dev_copies = std::vector>{}; - test_dev_copies.emplace_back(dev_buf, device_type::gpu); - test_dev_copies.emplace_back(dev_buf, device_type::gpu, 0); - test_dev_copies.emplace_back(dev_buf, device_type::gpu, 0, cuda_stream{}); - for (auto& copy_buf : test_dev_copies) { - data_out = std::vector(data.size()); - cuda_check(cudaMemcpy(static_cast(data_out.data()), - static_cast(copy_buf.data()), - copy_buf.size() * sizeof(int), - cudaMemcpyDefault)); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); - } - - auto test_host_buffers = std::vector>{}; - test_host_buffers.emplace_back(dev_buf, device_type::cpu); - test_host_buffers.emplace_back(dev_buf, device_type::cpu, 0); - test_host_buffers.emplace_back(dev_buf, device_type::cpu, 0, cuda_stream{}); - for (auto& host_buf : test_host_buffers) { - data_out = std::vector(host_buf.data(), host_buf.data() + host_buf.size()); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); - } - } -#endif - } -} - -TEST(Buffer, move_buffer) -{ - auto data = std::vector{1, 2, 3}; - auto test_buffers = std::vector>{}; - test_buffers.emplace_back(buffer(data.data(), data.size(), device_type::cpu)); - test_buffers.emplace_back(buffer(data.data(), data.size(), device_type::cpu), - device_type::cpu); - test_buffers.emplace_back( - buffer(data.data(), data.size(), device_type::cpu), device_type::cpu, 0); - test_buffers.emplace_back( - buffer(data.data(), data.size(), device_type::cpu), device_type::cpu, 0, cuda_stream{}); - - for (auto& buf : test_buffers) { - ASSERT_EQ(buf.memory_type(), device_type::cpu); - ASSERT_EQ(buf.size(), data.size()); - ASSERT_EQ(buf.data(), data.data()); - - auto data_out = std::vector(buf.data(), buf.data() + buf.size()); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); - } -#ifdef CUML_ENABLE_GPU - test_buffers = std::vector>{}; - test_buffers.emplace_back(buffer(data.data(), data.size(), device_type::cpu), - device_type::gpu); - test_buffers.emplace_back( - buffer(data.data(), data.size(), device_type::cpu), device_type::gpu, 0); - test_buffers.emplace_back( - buffer(data.data(), data.size(), device_type::cpu), device_type::gpu, 0, cuda_stream{}); - for (auto& buf : test_buffers) { - ASSERT_EQ(buf.memory_type(), device_type::gpu); - ASSERT_EQ(buf.size(), data.size()); - ASSERT_NE(buf.data(), data.data()); - - auto data_out = std::vector(buf.size()); - cuda_check(cudaMemcpy(static_cast(data_out.data()), - static_cast(buf.data()), - buf.size() * sizeof(int), - cudaMemcpyDefault)); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(data)); - } -#endif -} - -TEST(Buffer, move_assignment_buffer) -{ - auto data = std::vector{1, 2, 3}; - -#ifdef CUML_ENABLE_GPU - auto buf = buffer{data.data(), data.size() - 1, device_type::gpu}; -#else - auto buf = buffer{data.data(), data.size() - 1, device_type::cpu}; -#endif - buf = buffer{data.size(), device_type::cpu}; - - ASSERT_EQ(buf.memory_type(), device_type::cpu); - ASSERT_EQ(buf.size(), data.size()); -} - -TEST(Buffer, partial_buffer_copy) -{ - auto data1 = std::vector{1, 2, 3, 4, 5}; - auto data2 = std::vector{0, 0, 0, 0, 0}; - auto expected = std::vector{0, 3, 4, 5, 0}; -#ifdef CUML_ENABLE_GPU - auto buf1 = - buffer{buffer{data1.data(), data1.size(), device_type::cpu}, device_type::gpu}; -#else - auto buf1 = buffer{data1.data(), data1.size(), device_type::cpu}; -#endif - auto buf2 = buffer{data2.data(), data2.size(), device_type::cpu}; - copy(buf2, buf1, 1, 2, 3, cuda_stream{}); - copy(buf2, buf1, 1, 2, 3, cuda_stream{}); - EXPECT_THROW(copy(buf2, buf1, 1, 2, 4, cuda_stream{}), out_of_bounds); - EXPECT_THROW(copy(buf2, buf1, 1, data1.size() + 1, 0, cuda_stream{}), out_of_bounds); - EXPECT_THROW(copy(buf2, buf1, data2.size() + 1, 2, 0, cuda_stream{}), out_of_bounds); - EXPECT_THROW(copy(buffer{data2.data(), data2.size(), device_type::cpu}, - buffer{data1.data(), data1.size(), device_type::cpu}, - 1, - data1.size() + 1, - 0, - cuda_stream{}), - out_of_bounds); - EXPECT_THROW(copy(buffer{data2.data(), data2.size(), device_type::cpu}, - buffer{data1.data(), data1.size(), device_type::cpu}, - data2.size() + 1, - 2, - 0, - cuda_stream{}), - out_of_bounds); -} - -TEST(Buffer, buffer_copy_overloads) -{ - auto data = std::vector{1, 2, 3}; - auto expected = data; - auto orig_host_buffer = buffer(data.data(), data.size(), device_type::cpu); - auto orig_dev_buffer = buffer(orig_host_buffer, device_type::gpu); - auto copy_dev_buffer = buffer(data.size(), device_type::gpu); - - // copying host to host - auto data_out = std::vector(data.size()); - auto copy_host_buffer = buffer(data_out.data(), data.size(), device_type::cpu); - copy(copy_host_buffer, orig_host_buffer); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(expected)); - - // copying host to host with stream - data_out = std::vector(data.size()); - copy_host_buffer = buffer(data_out.data(), data.size(), device_type::cpu); - copy(copy_host_buffer, orig_host_buffer, cuda_stream{}); - EXPECT_THAT(data_out, ::testing::ElementsAreArray(expected)); - - // copying host to host with offset - data_out = std::vector(data.size() + 1); - copy_host_buffer = buffer(data_out.data(), data.size(), device_type::cpu); - copy(copy_host_buffer, orig_host_buffer, 2, 1, 1, cuda_stream{}); - expected = std::vector{0, 0, 2, 0}; - EXPECT_THAT(data_out, ::testing::ElementsAreArray(expected)); - -#ifdef CUML_ENABLE_GPU - // copy device to host - data_out = std::vector(data.size()); - copy_host_buffer = buffer(data_out.data(), data.size(), device_type::cpu); - copy(copy_host_buffer, orig_dev_buffer); - expected = data; - EXPECT_THAT(data_out, ::testing::ElementsAreArray(expected)); - - // copy device to host with stream - data_out = std::vector(data.size()); - copy_host_buffer = buffer(data_out.data(), data.size(), device_type::cpu); - copy(copy_host_buffer, orig_dev_buffer, cuda_stream{}); - expected = data; - EXPECT_THAT(data_out, ::testing::ElementsAreArray(expected)); - - // copy device to host with offset - data_out = std::vector(data.size() + 1); - copy_host_buffer = buffer(data_out.data(), data.size(), device_type::cpu); - copy(copy_host_buffer, orig_dev_buffer, 2, 1, 1, cuda_stream{}); - expected = std::vector{0, 0, 2, 0}; - EXPECT_THAT(data_out, ::testing::ElementsAreArray(expected)); -#endif -} - -} // namespace raft_proto diff --git a/cpp/tests/sg/fil/raft_proto/buffer.cu b/cpp/tests/sg/fil/raft_proto/buffer.cu deleted file mode 100644 index 16893688e2..0000000000 --- a/cpp/tests/sg/fil/raft_proto/buffer.cu +++ /dev/null @@ -1,42 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include -#include -#include -#include -#include - -#include - -#include -#include - -#include - -namespace raft_proto { - -CUML_KERNEL void check_buffer_access(int* buf) -{ - if (buf[0] == 1) { buf[0] = 4; } - if (buf[1] == 2) { buf[1] = 5; } - if (buf[2] == 3) { buf[2] = 6; } -} - -TEST(Buffer, device_buffer_access) -{ - auto data = std::vector{1, 2, 3}; - auto expected = std::vector{4, 5, 6}; - auto buf = buffer( - buffer(data.data(), data.size(), device_type::cpu), device_type::gpu, 0, cuda_stream{}); - check_buffer_access<<<1, 1>>>(buf.data()); - auto data_out = std::vector(expected.size()); - auto host_buf = buffer(data_out.data(), data_out.size(), device_type::cpu); - copy(host_buf, buf); - ASSERT_EQ(cudaStreamSynchronize(cuda_stream{}), cudaSuccess); - EXPECT_THAT(data_out, testing::ElementsAreArray(expected)); -} - -} // namespace raft_proto diff --git a/cpp/tests/sg/fil/treelite_importer.cpp b/cpp/tests/sg/fil/treelite_importer.cpp deleted file mode 100644 index 6ab4d45b59..0000000000 --- a/cpp/tests/sg/fil/treelite_importer.cpp +++ /dev/null @@ -1,378 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include -#include -#include -#include - -#include -#include -#include -#include -#include - -#include -#include -#include - -namespace ML { -namespace fil { - -/* |Test Forest - * |-----------------------------------|--------------------------------------| - * |TREE 0 |KEY | - * | A-0-3 |* Non-leaf nodes: Label-Fea-Split | - * | / \ |* Cat nodes: Label-Fea-{cat0, cat1...}| - * | / \ |* Leaf nodes: Label above outputs | - * | / \ | - Regression output | - * | / \ | - Binary classification output | - * | B-1-2 C | - Multiclass output | - * | /\ 0 |--------------------------------------| - * | / \ | - * | D E-3-0 | - * | 1 /\ | - * | F G | - * | 2 3 | - * |-----------------------------------| - * |TREE 1 | - * | H-0-0 | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | I J-1-1 | - * | 4 /\ | - * | / \ | - * | K L | - * | 5 6 | - * |-----------------------------------| - * |TREE 2 | - * | M-5-{0, 1, 3} | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | N O | - * | 7 8 | - * |-----------------------------------| - * |TREE 3 | - * | P | - * | 9 | - * | | - * | | - * | | - * | | - * |-----------------------------------| - * |TREE 4 | - * | Q-5-{0, 7} | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | R S | - * | 10 11 | - * |-----------------------------------| - * |TREE 5 | - * | T-6-{1, 3} | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | U V-6-{1, 4} | - * | 12 /\ | - * | / \ | - * | W-6-{3} X | - * | /\ 13 | - * | Y Z | - * | 14 15 | - * |-----------------------------------| - */ -auto static constexpr const SAMPLE_COL_COUNT = 7; -auto static constexpr const SAMPLE_TREE_COUNT = 6; -auto static const SAMPLE_FOREST = []() { - auto metadata = treelite::model_builder::Metadata{ - SAMPLE_COL_COUNT, - treelite::TaskType::kRegressor, - true, - 1, - {1}, - {1, 1}, - }; - auto tree_annotation = - treelite::model_builder::TreeAnnotation{SAMPLE_TREE_COUNT, - std::vector(SAMPLE_TREE_COUNT, 0), - std::vector{0, 0, 0, 0, 0, 0}}; - auto model_builder = treelite::model_builder::GetModelBuilder( - treelite::TypeInfo::kFloat32, - treelite::TypeInfo::kFloat32, - metadata, - tree_annotation, - treelite::model_builder::PostProcessorFunc{"identity_multiclass"}, - std::vector(1, 0.0f)); - // TREE 0 - model_builder->StartTree(); - // Node A - model_builder->StartNode(0); - // For numerical splits, the right child is "hot" if the operator is kLT or - // kLE. For categorical splits, whichever child corresponds to - // out-of-category is the "hot" node. - // feature index, threshold, default left, operator, left child, right child - model_builder->NumericalTest(0, 3.0, true, treelite::Operator::kGE, 1, 2); - model_builder->EndNode(); - // Node B - model_builder->StartNode(1); - model_builder->NumericalTest(1, 2.0, false, treelite::Operator::kLT, 4, 3); - model_builder->EndNode(); - // Node C - model_builder->StartNode(2); - model_builder->LeafScalar(0.0); - model_builder->EndNode(); - // Node D - model_builder->StartNode(3); - model_builder->LeafScalar(1.0); - model_builder->EndNode(); - // Node E - model_builder->StartNode(4); - model_builder->NumericalTest(3, 0.0, true, treelite::Operator::kGT, 5, 6); - model_builder->EndNode(); - // Node F - model_builder->StartNode(5); - model_builder->LeafScalar(2.0); - model_builder->EndNode(); - // Node G - model_builder->StartNode(6); - model_builder->LeafScalar(3.0); - model_builder->EndNode(); - model_builder->EndTree(); - - // TREE 1 - model_builder->StartTree(); - // Node H - model_builder->StartNode(0); - model_builder->NumericalTest(0, 0.0, true, treelite::Operator::kGE, 1, 2); - model_builder->EndNode(); - // Node I - model_builder->StartNode(1); - model_builder->LeafScalar(4.0); - model_builder->EndNode(); - // Node J - model_builder->StartNode(2); - model_builder->NumericalTest(1, 0.0, true, treelite::Operator::kGE, 3, 4); - model_builder->EndNode(); - // Node K - model_builder->StartNode(3); - model_builder->LeafScalar(5.0); - model_builder->EndNode(); - // Node L - model_builder->StartNode(4); - model_builder->LeafScalar(6.0); - model_builder->EndNode(); - model_builder->EndTree(); - - // TREE 2 - model_builder->StartTree(); - // Node M - model_builder->StartNode(0); - model_builder->CategoricalTest(5, true, std::vector{0, 1, 3}, true, 1, 2); - model_builder->EndNode(); - // Node N - model_builder->StartNode(1); - model_builder->LeafScalar(7.0); - model_builder->EndNode(); - // Node O - model_builder->StartNode(2); - model_builder->LeafScalar(8.0); - model_builder->EndNode(); - model_builder->EndTree(); - - // TREE 3 - model_builder->StartTree(); - // Node P - model_builder->StartNode(0); - model_builder->LeafScalar(9.0); - model_builder->EndNode(); - model_builder->EndTree(); - - // TREE 4 - model_builder->StartTree(); - // Node Q - model_builder->StartNode(0); - model_builder->CategoricalTest(5, true, std::vector{0, 7}, false, 2, 1); - model_builder->EndNode(); - // Node R - model_builder->StartNode(1); - model_builder->LeafScalar(10.0); - model_builder->EndNode(); - // Node S - model_builder->StartNode(2); - model_builder->LeafScalar(11.0); - model_builder->EndNode(); - model_builder->EndTree(); - - // TREE 5 - model_builder->StartTree(); - // Node T - model_builder->StartNode(0); - model_builder->CategoricalTest(6, true, std::vector{1, 3}, true, 1, 2); - model_builder->EndNode(); - // Node U - model_builder->StartNode(1); - model_builder->LeafScalar(12.0); - model_builder->EndNode(); - // Node V - model_builder->StartNode(2); - model_builder->CategoricalTest(6, true, std::vector{1, 4}, true, 3, 4); - model_builder->EndNode(); - // Node W - model_builder->StartNode(3); - model_builder->CategoricalTest(6, true, std::vector{3}, true, 5, 6); - model_builder->EndNode(); - // Node X - model_builder->StartNode(4); - model_builder->LeafScalar(13.0); - model_builder->EndNode(); - // Node Y - model_builder->StartNode(5); - model_builder->LeafScalar(14.0); - model_builder->EndNode(); - // Node Z - model_builder->StartNode(6); - model_builder->LeafScalar(15.0); - model_builder->EndNode(); - model_builder->EndTree(); - return model_builder->CommitModel(); -}(); - -TEST(TreeliteImporter, depth_first) -{ - auto fil_model = import_from_treelite_model(*SAMPLE_FOREST, tree_layout::depth_first); - ASSERT_EQ(fil_model.num_features(), 7); - ASSERT_EQ(fil_model.num_outputs(), 1); - ASSERT_EQ(fil_model.num_trees(), 6); - ASSERT_FALSE(fil_model.has_vector_leaves()); - ASSERT_EQ(fil_model.row_postprocessing(), row_op::disable); - ASSERT_EQ(fil_model.elem_postprocessing(), element_op::disable); - ASSERT_EQ(fil_model.memory_type(), raft_proto::device_type::cpu); - ASSERT_EQ(fil_model.device_index(), 0); - ASSERT_FALSE(fil_model.is_double_precision()); -} - -TEST(TreeliteImporter, breadth_first) -{ - auto fil_model = import_from_treelite_model(*SAMPLE_FOREST, tree_layout::breadth_first); - ASSERT_EQ(fil_model.num_features(), 7); - ASSERT_EQ(fil_model.num_outputs(), 1); - ASSERT_EQ(fil_model.num_trees(), 6); - ASSERT_FALSE(fil_model.has_vector_leaves()); - ASSERT_EQ(fil_model.row_postprocessing(), row_op::disable); - ASSERT_EQ(fil_model.elem_postprocessing(), element_op::disable); - ASSERT_EQ(fil_model.memory_type(), raft_proto::device_type::cpu); - ASSERT_EQ(fil_model.device_index(), 0); - ASSERT_FALSE(fil_model.is_double_precision()); -} - -TEST(TreeliteImporter, layered_children_together) -{ - auto fil_model = - import_from_treelite_model(*SAMPLE_FOREST, tree_layout::layered_children_together); - ASSERT_EQ(fil_model.num_features(), 7); - ASSERT_EQ(fil_model.num_outputs(), 1); - ASSERT_EQ(fil_model.num_trees(), 6); - ASSERT_FALSE(fil_model.has_vector_leaves()); - ASSERT_EQ(fil_model.row_postprocessing(), row_op::disable); - ASSERT_EQ(fil_model.elem_postprocessing(), element_op::disable); - ASSERT_EQ(fil_model.memory_type(), raft_proto::device_type::cpu); - ASSERT_EQ(fil_model.device_index(), 0); - ASSERT_FALSE(fil_model.is_double_precision()); -} - -template -auto make_degenerate_tree(const leaf_t& leaf) -{ - auto task_type = treelite::TaskType{}; - auto num_class = std::int32_t{}; - auto class_annotation = std::vector{}; - if constexpr (use_leaf_vector) { - task_type = treelite::TaskType::kMultiClf; - num_class = leaf.size(); - class_annotation = {-1}; - } else { - task_type = treelite::TaskType::kBinaryClf; - num_class = 1; - class_annotation = {0}; - } - auto metadata = treelite::model_builder::Metadata{ - 1, - task_type, - false, - 1, - {num_class}, - {1, num_class}, - }; - auto tree_annotation = treelite::model_builder::TreeAnnotation{1, {0}, class_annotation}; - auto model_builder = treelite::model_builder::GetModelBuilder( - treelite::TypeInfo::kFloat64, - treelite::TypeInfo::kFloat64, - metadata, - tree_annotation, - treelite::model_builder::PostProcessorFunc{"identity_multiclass"}, - std::vector(num_class, 0.0)); - model_builder->StartTree(); - model_builder->StartNode(0); - if constexpr (use_leaf_vector) { - model_builder->LeafVector(leaf); - } else { - model_builder->LeafScalar(leaf); - } - model_builder->EndNode(); - model_builder->EndTree(); - return model_builder->CommitModel(); -} - -TEST(TreeliteImporter, DegenerateTree) -{ - auto tl_model = make_degenerate_tree(1.0); - auto fil_model = import_from_treelite_model(*tl_model, tree_layout::breadth_first); - ASSERT_FALSE(fil_model.has_vector_leaves()); - - auto handle = raft::handle_t{}; - auto X = std::vector{0.0}; - auto preds = std::vector(1, 0.0); - auto expected_preds = std::vector{1.0}; - fil_model.predict(handle, - preds.data(), - X.data(), - 1, - raft_proto::device_type::cpu, - raft_proto::device_type::cpu, - ML::fil::infer_kind::default_kind, - 1); - ASSERT_EQ(preds, expected_preds); -} - -TEST(TreeliteImporter, DegenerateTreeWithVectorLeaf) -{ - auto tl_model = make_degenerate_tree(std::vector{0.5, 0.5}); - auto fil_model = import_from_treelite_model(*tl_model, tree_layout::breadth_first); - ASSERT_TRUE(fil_model.has_vector_leaves()); - - auto handle = raft::handle_t{}; - auto X = std::vector{0.0}; - auto preds = std::vector(2, 0.0); - auto expected_preds = std::vector{0.5, 0.5}; - fil_model.predict(handle, - preds.data(), - X.data(), - 1, - raft_proto::device_type::cpu, - raft_proto::device_type::cpu, - ML::fil::infer_kind::default_kind, - 1); - ASSERT_EQ(preds, expected_preds); -} - -} // namespace fil -} // namespace ML diff --git a/cpp/tests/sg/fil/treelite_importer_invalid_inputs.cpp b/cpp/tests/sg/fil/treelite_importer_invalid_inputs.cpp deleted file mode 100644 index aa0bae0661..0000000000 --- a/cpp/tests/sg/fil/treelite_importer_invalid_inputs.cpp +++ /dev/null @@ -1,248 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include -#include -#include -#include - -#include -#include -#include -#include -#include - -#include -#include -#include - -namespace ML { -namespace fil { - -TEST(TreeliteImporter, large_category_value) -{ - // For tree models with 32-bit data storage, - // attempting to use UINT32_MAX in a categorical test node - // must throw an exception. - - auto metadata = treelite::model_builder::Metadata{ - 1, - treelite::TaskType::kRegressor, - false, - 1, - {1}, - {1, 1}, - }; - auto tree_annotation = treelite::model_builder::TreeAnnotation{1, {0}, {0}}; - auto model_builder = - treelite::model_builder::GetModelBuilder(treelite::TypeInfo::kFloat32, - treelite::TypeInfo::kFloat32, - metadata, - tree_annotation, - treelite::model_builder::PostProcessorFunc{"identity"}, - {0.0}); - - model_builder->StartTree(); - model_builder->StartNode(0); - model_builder->CategoricalTest( - 0, false, {std::numeric_limits::max()}, false, 1, 2); - model_builder->EndNode(); - - model_builder->StartNode(1); - model_builder->LeafScalar(1.0f); - model_builder->EndNode(); - - model_builder->StartNode(2); - model_builder->LeafScalar(-1.0f); - model_builder->EndNode(); - - model_builder->EndTree(); - - auto tl_model = model_builder->CommitModel(); - - auto expected_error_msg = std::string{"Tree 0, Node 0: Category index must be at most "} + - std::to_string(std::numeric_limits::max() - 1); - - ASSERT_THAT([&]() { import_from_treelite_model(*tl_model, tree_layout::breadth_first); }, - testing::ThrowsMessage(testing::HasSubstr(expected_error_msg))); -} - -TEST(TreeliteImporter, large_category_value2) -{ - // For tree models with 32-bit data storage, - // it should be possible to use (UINT32_MAX - 1) in a categorical test node - - auto metadata = treelite::model_builder::Metadata{ - 1, - treelite::TaskType::kRegressor, - false, - 1, - {1}, - {1, 1}, - }; - auto tree_annotation = treelite::model_builder::TreeAnnotation{1, {0}, {0}}; - auto model_builder = - treelite::model_builder::GetModelBuilder(treelite::TypeInfo::kFloat32, - treelite::TypeInfo::kFloat32, - metadata, - tree_annotation, - treelite::model_builder::PostProcessorFunc{"identity"}, - {0.0}); - - model_builder->StartTree(); - model_builder->StartNode(0); - model_builder->CategoricalTest( - 0, false, {std::numeric_limits::max() - 1}, false, 1, 2); - model_builder->EndNode(); - - model_builder->StartNode(1); - model_builder->LeafScalar(1.0f); - model_builder->EndNode(); - - model_builder->StartNode(2); - model_builder->LeafScalar(-1.0f); - model_builder->EndNode(); - - model_builder->EndTree(); - - auto tl_model = model_builder->CommitModel(); - ASSERT_NO_THROW(import_from_treelite_model(*tl_model, tree_layout::breadth_first)); -} - -TEST(TreeliteImporter, large_feature_id) -{ - // Tree models with 16-bit storage for node metadata should throw - // an exception for feature IDs larger than 0x1FFF. - - auto metadata = treelite::model_builder::Metadata{ - 9000, - treelite::TaskType::kRegressor, - false, - 1, - {1}, - {1, 1}, - }; - auto tree_annotation = treelite::model_builder::TreeAnnotation{1, {0}, {0}}; - auto model_builder = - treelite::model_builder::GetModelBuilder(treelite::TypeInfo::kFloat32, - treelite::TypeInfo::kFloat32, - metadata, - tree_annotation, - treelite::model_builder::PostProcessorFunc{"identity"}, - {0.0}); - - model_builder->StartTree(); - model_builder->StartNode(0); - // Use a "large" feature ID here - model_builder->NumericalTest(8999, 0.0, false, treelite::Operator::kGT, 1, 2); - model_builder->EndNode(); - - model_builder->StartNode(1); - model_builder->LeafScalar(1.0f); - model_builder->EndNode(); - - model_builder->StartNode(2); - model_builder->LeafScalar(-1.0f); - model_builder->EndNode(); - - model_builder->EndTree(); - - auto tl_model = model_builder->CommitModel(); - - // Normally, treelite_importer::import() would choose the right size - // for the metadata storage, sufficient to hold all given feature IDs. - // For this example, it chooses 32-bit storage type (due to the use of feature ID 8999). - ASSERT_NO_THROW(import_from_treelite_model(*tl_model, tree_layout::breadth_first)); - - // Trick the importer to pick 16-bit storage type for metadata storage. - auto variant_index = get_forest_variant_index(false, 2, 1); - auto importer = treelite_importer{}; - - // The importer should throw an informative error message rather than silently - // truncating the feature ID. - auto expected_error_msg = - std::string{"Tree 0, Node 0: The 'feature' value in the node must be at most "} + - std::to_string(0x1FFF); - ASSERT_THAT( - [&] { - importer.import_to_specific_variant(variant_index, - *tl_model, - importer.get_num_class(*tl_model), - importer.get_num_feature(*tl_model), - importer.get_max_num_categories(*tl_model), - importer.get_offsets(*tl_model)); - }, - testing::ThrowsMessage(testing::HasSubstr(expected_error_msg))); -} - -TEST(TreeliteImporter, safe_cast_floating_point) -{ - /* Valid casts */ - ASSERT_NO_THROW( - detail::safe_cast_floating_point(double{3.1})); // Some loss of precision, but o.k. - ASSERT_NO_THROW(detail::safe_cast_floating_point(std::numeric_limits::max())); - - // INFs and NANs are allowed for widening cast - ASSERT_NO_THROW(detail::safe_cast_floating_point(std::numeric_limits::infinity())); - ASSERT_NO_THROW( - detail::safe_cast_floating_point(std::numeric_limits::infinity())); - ASSERT_NO_THROW(detail::safe_cast_floating_point(std::numeric_limits::infinity())); - ASSERT_NO_THROW(detail::safe_cast_floating_point(std::numeric_limits::quiet_NaN())); - ASSERT_NO_THROW( - detail::safe_cast_floating_point(std::numeric_limits::quiet_NaN())); - ASSERT_NO_THROW( - detail::safe_cast_floating_point(std::numeric_limits::quiet_NaN())); - - // Invalid casts - auto inf_msg = std::string{"Cannot cast an INF or NaN value"}; - ASSERT_THAT( - [] { detail::safe_cast_floating_point(std::numeric_limits::infinity()); }, - testing::ThrowsMessage(testing::HasSubstr(inf_msg))); - ASSERT_THAT( - [] { detail::safe_cast_floating_point(std::numeric_limits::quiet_NaN()); }, - testing::ThrowsMessage(testing::HasSubstr(inf_msg))); - ASSERT_THAT([] { detail::safe_cast_floating_point(double{1e100}); }, - testing::ThrowsMessage( - testing::HasSubstr("Input must be at most"))); - ASSERT_THAT([] { detail::safe_cast_floating_point(double{-1e100}); }, - testing::ThrowsMessage( - testing::HasSubstr("Input must be at least"))); -} - -TEST(TreeliteImporter, invalid_postproc_constant) -{ - auto metadata = treelite::model_builder::Metadata{ - 1, - treelite::TaskType::kRegressor, - false, - 1, - {1}, - {1, 1}, - }; - auto tree_annotation = treelite::model_builder::TreeAnnotation{1, {0}, {0}}; - auto model_builder = treelite::model_builder::GetModelBuilder( - treelite::TypeInfo::kFloat32, - treelite::TypeInfo::kFloat32, - metadata, - tree_annotation, - treelite::model_builder::PostProcessorFunc{ - "sigmoid", {{"sigmoid_alpha", std::numeric_limits::quiet_NaN()}}}, - {0.0}); - - model_builder->StartTree(); - model_builder->StartNode(0); - model_builder->LeafScalar(0.0f); - model_builder->EndNode(); - model_builder->EndTree(); - - auto tl_model = model_builder->CommitModel(); - - ASSERT_THAT([&] { import_from_treelite_model(*tl_model, tree_layout::breadth_first); }, - testing::ThrowsMessage( - testing::HasSubstr("Found an invalid value for postprocessing constant"))); -} - -} // namespace fil -} // namespace ML diff --git a/cpp/tests/sg/forest/traversal_forest.cpp b/cpp/tests/sg/forest/traversal_forest.cpp deleted file mode 100644 index 4ddd333324..0000000000 --- a/cpp/tests/sg/forest/traversal_forest.cpp +++ /dev/null @@ -1,327 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include -#include -#include - -#include -#include - -#include - -namespace ML { -namespace forest { - -struct test_node : traversal_node<> { - test_node() : label_{}, hot_child_{}, distant_child_{} {} - test_node(char label, - std::optional hot_child, - std::optional distant_child) - : label_{label}, hot_child_{hot_child}, distant_child_{distant_child} - { - } - auto get_label() { return label_; } - bool is_leaf() const override { return !hot_child_.has_value(); } - std::size_t hot_child() const override { return hot_child_.value_or(std::size_t{}); } - std::size_t distant_child() const override { return distant_child_.value_or(std::size_t{}); } - - private: - char label_; - std::optional hot_child_{}; - std::optional distant_child_{}; -}; - -/* |Test Forest - * |-----------------------------------| - * |TREE 0 | - * | A | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | B C | - * | /\ | - * | / \ | - * | D E | - * | /\ | - * | F G | - * |-----------------------------------| - * |TREE 1 | - * | H | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | I J | - * | /\ | - * | / \ | - * | K L | - * |-----------------------------------| - * |TREE 2 | - * | M | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | N O | - * |-----------------------------------| - * |TREE 3 | - * | P | - * | | - * | | - * | | - * | | - * | | - * |-----------------------------------| - * |TREE 4 | - * | Q | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | R S | - * |-----------------------------------| - * |TREE 5 | - * | T | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | U V | - * | /\ | - * | / \ | - * | W X | - * | /\ | - * | Y Z | - * |-----------------------------------| - */ - -struct test_forest : traversal_forest { - test_forest() - : traversal_forest{std::vector>{ - std::make_pair(std::size_t{}, std::size_t{}), - std::make_pair(std::size_t{1}, std::size_t{7}), - std::make_pair(std::size_t{2}, std::size_t{12}), - std::make_pair(std::size_t{3}, std::size_t{15}), - std::make_pair(std::size_t{4}, std::size_t{16}), - std::make_pair(std::size_t{5}, std::size_t{19}), - }} - { - } - test_node get_node(std::size_t tree_id, std::size_t node_id) const override - { - return nodes_[node_id]; - } - - private: - std::vector nodes_{ - test_node{'A', std::size_t{1}, std::size_t{2}}, // 0 - test_node{'B', std::size_t{3}, std::size_t{4}}, // 1 - test_node{'C', std::optional{}, std::optional{}}, // 2 - test_node{'D', std::optional{}, std::optional{}}, // 3 - test_node{'E', std::size_t{5}, std::size_t{6}}, // 4 - test_node{'F', std::optional{}, std::optional{}}, // 5 - test_node{'G', std::optional{}, std::optional{}}, // 6 - test_node{'H', std::size_t{8}, std::size_t{9}}, // 7 - test_node{'I', std::optional{}, std::optional{}}, // 8 - test_node{'J', std::size_t{10}, std::size_t{11}}, // 9 - test_node{'K', std::optional{}, std::optional{}}, // 10 - test_node{'L', std::optional{}, std::optional{}}, // 11 - test_node{'M', std::size_t{13}, std::size_t{14}}, // 12 - test_node{'N', std::optional{}, std::optional{}}, // 13 - test_node{'O', std::optional{}, std::optional{}}, // 14 - test_node{'P', std::optional{}, std::optional{}}, // 15 - test_node{'Q', std::size_t{17}, std::size_t{18}}, // 16 - test_node{'R', std::optional{}, std::optional{}}, // 17 - test_node{'S', std::optional{}, std::optional{}}, // 18 - test_node{'T', std::size_t{20}, std::size_t{21}}, // 19 - test_node{'U', std::optional{}, std::optional{}}, // 20 - test_node{'V', std::size_t{22}, std::size_t{23}}, // 21 - test_node{'W', std::size_t{24}, std::size_t{25}}, // 22 - test_node{'X', std::optional{}, std::optional{}}, // 23 - test_node{'Y', std::optional{}, std::optional{}}, // 24 - test_node{'Z', std::optional{}, std::optional{}} // 25 - }; -}; - -struct traversal_forest_results { - std::string order; - std::vector depth; - std::vector parents; - std::vector tree_indices; -}; - -auto static const TRAVERSAL_RESULTS = - std::vector>{ - std::make_pair(forest_order::depth_first, - traversal_forest_results{ - "ABDEFGCHIJKLMNOPQRSTUVWYZX", - std::vector{0, 1, 2, 2, 3, 3, 1, 0, 1, 1, 2, 2, 0, - 1, 1, 0, 0, 1, 1, 0, 1, 1, 2, 3, 3, 2}, - std::vector{0, 0, 1, 1, 3, 3, 0, 7, 7, 7, 9, 9, 12, - 12, 12, 15, 16, 16, 16, 19, 19, 19, 21, 22, 22, 21}, - std::vector{0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 2, - 2, 2, 3, 4, 4, 4, 5, 5, 5, 5, 5, 5, 5}}), - std::make_pair(forest_order::breadth_first, - traversal_forest_results{ - "ABCDEFGHIJKLMNOPQRSTUVWXYZ", - std::vector{0, 1, 1, 2, 2, 3, 3, 0, 1, 1, 2, 2, 0, - 1, 1, 0, 0, 1, 1, 0, 1, 1, 2, 2, 3, 3}, - std::vector{0, 0, 0, 1, 1, 4, 4, 7, 7, 7, 9, 9, 12, - 12, 12, 15, 16, 16, 16, 19, 19, 19, 21, 21, 22, 22}, - std::vector{0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 2, - 2, 2, 3, 4, 4, 4, 5, 5, 5, 5, 5, 5, 5}}), - std::make_pair(forest_order::layered_children_together, - traversal_forest_results{ - "AHMPQTBCIJNORSUVDEKLWXFGYZ", - std::vector{0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 2, 2, 2, 2, 2, 2, 3, 3, 3, 3}, - std::vector{0, 1, 2, 3, 4, 5, 0, 0, 1, 1, 2, 2, 4, - 4, 5, 5, 6, 6, 9, 9, 15, 15, 17, 17, 20, 20}, - std::vector{0, 1, 2, 3, 4, 5, 0, 0, 1, 1, 2, 2, 4, - 4, 5, 5, 0, 0, 1, 1, 5, 5, 0, 0, 5, 5}}), - std::make_pair(forest_order::layered_children_segregated, - traversal_forest_results{ - "AHMPQTBINRUCJOSVDKWELXYFZG", - std::vector{0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 2, 2, 2, 2, 2, 2, 3, 3, 3, 3}, - std::vector{0, 1, 2, 3, 4, 5, 0, 1, 2, 4, 5, 0, 1, - 2, 4, 5, 6, 12, 15, 6, 12, 15, 18, 19, 18, 19}, - std::vector{0, 1, 2, 3, 4, 5, 0, 1, 2, 4, 5, 0, 1, - 2, 4, 5, 0, 1, 5, 0, 1, 5, 5, 0, 5, 0}}), - }; - -template -auto get_expected_for_each_result() -{ - return std::find_if(std::begin(TRAVERSAL_RESULTS), - std::end(TRAVERSAL_RESULTS), - [](auto&& pair) { return pair.first == order; }) - ->second; -} - -template -auto get_for_each_order() -{ - auto result = std::vector{}; - test_forest{}.for_each( - [&result](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - result.push_back(node.get_label()); - }); - return std::string(std::begin(result), std::end(result)); -} - -template -auto get_for_each_depth() -{ - auto result = std::vector{}; - test_forest{}.for_each( - [&result](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - result.push_back(depth); - }); - return result; -} - -template -auto get_for_each_parent() -{ - auto result = std::vector{}; - test_forest{}.for_each( - [&result](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - result.push_back(parent_index); - }); - return result; -} - -template -auto get_for_each_tree() -{ - auto result = std::vector{}; - test_forest{}.for_each( - [&result](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - result.push_back(tree_id); - }); - return result; -} - -TEST(ForestTraversal, depth_first) -{ - auto order = get_for_each_order(); - auto depths = get_for_each_depth(); - auto parents = get_for_each_parent(); - auto trees = get_for_each_tree(); - auto expected = get_expected_for_each_result(); - EXPECT_EQ(order, expected.order); - for (auto i = std::size_t{}; i < expected.depth.size(); ++i) { - EXPECT_EQ(depths[i], expected.depth[i]); - } - for (auto i = std::size_t{}; i < expected.parents.size(); ++i) { - EXPECT_EQ(parents[i], expected.parents[i]); - } - for (auto i = std::size_t{}; i < expected.tree_indices.size(); ++i) { - EXPECT_EQ(trees[i], expected.tree_indices[i]); - } -} - -TEST(ForestTraversal, breadth_first) -{ - auto order = get_for_each_order(); - auto depths = get_for_each_depth(); - auto parents = get_for_each_parent(); - auto trees = get_for_each_tree(); - auto expected = get_expected_for_each_result(); - EXPECT_EQ(order, expected.order); - for (auto i = std::size_t{}; i < expected.depth.size(); ++i) { - EXPECT_EQ(depths[i], expected.depth[i]); - } - for (auto i = std::size_t{}; i < expected.parents.size(); ++i) { - EXPECT_EQ(parents[i], expected.parents[i]); - } - for (auto i = std::size_t{}; i < expected.tree_indices.size(); ++i) { - EXPECT_EQ(trees[i], expected.tree_indices[i]); - } -} - -TEST(ForestTraversal, layered_children_together) -{ - auto order = get_for_each_order(); - auto depths = get_for_each_depth(); - auto parents = get_for_each_parent(); - auto trees = get_for_each_tree(); - auto expected = get_expected_for_each_result(); - EXPECT_EQ(order, expected.order); - for (auto i = std::size_t{}; i < expected.depth.size(); ++i) { - EXPECT_EQ(depths[i], expected.depth[i]); - } - for (auto i = std::size_t{}; i < expected.parents.size(); ++i) { - EXPECT_EQ(parents[i], expected.parents[i]); - } - for (auto i = std::size_t{}; i < expected.tree_indices.size(); ++i) { - EXPECT_EQ(trees[i], expected.tree_indices[i]); - } -} - -TEST(ForestTraversal, layered_children_segregated) -{ - auto order = get_for_each_order(); - auto depths = get_for_each_depth(); - auto parents = get_for_each_parent(); - auto trees = get_for_each_tree(); - auto expected = get_expected_for_each_result(); - EXPECT_EQ(order, expected.order); - for (auto i = std::size_t{}; i < expected.depth.size(); ++i) { - EXPECT_EQ(depths[i], expected.depth[i]); - } - for (auto i = std::size_t{}; i < expected.parents.size(); ++i) { - EXPECT_EQ(parents[i], expected.parents[i]); - } - for (auto i = std::size_t{}; i < expected.tree_indices.size(); ++i) { - EXPECT_EQ(trees[i], expected.tree_indices[i]); - } -} - -} // namespace forest -} // namespace ML diff --git a/cpp/tests/sg/forest/treelite_traversal.cpp b/cpp/tests/sg/forest/treelite_traversal.cpp deleted file mode 100644 index a8991e8bb9..0000000000 --- a/cpp/tests/sg/forest/treelite_traversal.cpp +++ /dev/null @@ -1,487 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include -#include - -#include -#include -#include -#include -#include - -#include - -namespace ML { -namespace forest { - -/* |Test Forest - * |-----------------------------------|--------------------------------------| - * |TREE 0 |KEY | - * | A-0-3 |* Non-leaf nodes: Label-Fea-Split | - * | / \ |* Cat nodes: Label-Fea-{cat0, cat1...}| - * | / \ |* Leaf nodes: Label above outputs | - * | / \ | - Regression output | - * | / \ | - Binary classification output | - * | B-1-2 C | - Multiclass output | - * | /\ 0 |--------------------------------------| - * | / \ | - * | D E-3-0 | - * | 1 /\ | - * | F G | - * | 2 3 | - * |-----------------------------------| - * |TREE 1 | - * | H-0-0 | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | I J-1-1 | - * | 4 /\ | - * | / \ | - * | K L | - * | 5 6 | - * |-----------------------------------| - * |TREE 2 | - * | M-5-{0, 1, 3} | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | N O | - * | 7 8 | - * |-----------------------------------| - * |TREE 3 | - * | P | - * | 9 | - * | | - * | | - * | | - * | | - * |-----------------------------------| - * |TREE 4 | - * | Q-5-{0, 7} | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | R S | - * | 10 11 | - * |-----------------------------------| - * |TREE 5 | - * | T-6-{1, 3} | - * | / \ | - * | / \ | - * | / \ | - * | / \ | - * | U V-6-{1, 4} | - * | 12 /\ | - * | / \ | - * | W-6-{3} X | - * | /\ 13 | - * | Y Z | - * | 14 15 | - * |-----------------------------------| - */ -auto static constexpr const SAMPLE_COL_COUNT = 7; -auto static constexpr const SAMPLE_TREE_COUNT = 6; -auto static constexpr const SAMPLE_CATEGORICAL_COUNT = 5; -auto static const SAMPLE_FOREST = []() { - auto metadata = treelite::model_builder::Metadata{ - SAMPLE_COL_COUNT, - treelite::TaskType::kRegressor, - true, - 1, - {1}, - {1, 1}, - }; - auto tree_annotation = - treelite::model_builder::TreeAnnotation{SAMPLE_TREE_COUNT, - std::vector(SAMPLE_TREE_COUNT, 0), - std::vector{0, 0, 0, 0, 0, 0}}; - auto model_builder = treelite::model_builder::GetModelBuilder( - treelite::TypeInfo::kFloat32, - treelite::TypeInfo::kFloat32, - metadata, - tree_annotation, - treelite::model_builder::PostProcessorFunc{"identity_multiclass"}, - std::vector(1, 0.0f)); - // TREE 0 - model_builder->StartTree(); - // Node A - model_builder->StartNode(0); - // For numerical splits, the right child is "hot" if the operator is kLT or - // kLE. For categorical splits, whichever child corresponds to - // out-of-category is the "hot" node. - // feature index, threshold, default left, operator, left child, right child - model_builder->NumericalTest(0, 3.0, true, treelite::Operator::kGE, 1, 2); - model_builder->EndNode(); - // Node B - model_builder->StartNode(1); - model_builder->NumericalTest(1, 2.0, false, treelite::Operator::kLT, 4, 3); - model_builder->EndNode(); - // Node C - model_builder->StartNode(2); - model_builder->LeafScalar(0.0); - model_builder->EndNode(); - // Node D - model_builder->StartNode(3); - model_builder->LeafScalar(1.0); - model_builder->EndNode(); - // Node E - model_builder->StartNode(4); - model_builder->NumericalTest(3, 0.0, true, treelite::Operator::kGT, 5, 6); - model_builder->EndNode(); - // Node F - model_builder->StartNode(5); - model_builder->LeafScalar(2.0); - model_builder->EndNode(); - // Node G - model_builder->StartNode(6); - model_builder->LeafScalar(3.0); - model_builder->EndNode(); - model_builder->EndTree(); - - // TREE 1 - model_builder->StartTree(); - // Node H - model_builder->StartNode(0); - model_builder->NumericalTest(0, 0.0, true, treelite::Operator::kGE, 1, 2); - model_builder->EndNode(); - // Node I - model_builder->StartNode(1); - model_builder->LeafScalar(4.0); - model_builder->EndNode(); - // Node J - model_builder->StartNode(2); - model_builder->NumericalTest(1, 0.0, true, treelite::Operator::kGE, 3, 4); - model_builder->EndNode(); - // Node K - model_builder->StartNode(3); - model_builder->LeafScalar(5.0); - model_builder->EndNode(); - // Node L - model_builder->StartNode(4); - model_builder->LeafScalar(6.0); - model_builder->EndNode(); - model_builder->EndTree(); - - // TREE 2 - model_builder->StartTree(); - // Node M - model_builder->StartNode(0); - model_builder->CategoricalTest(5, true, std::vector{0, 1, 3}, true, 1, 2); - model_builder->EndNode(); - // Node N - model_builder->StartNode(1); - model_builder->LeafScalar(7.0); - model_builder->EndNode(); - // Node O - model_builder->StartNode(2); - model_builder->LeafScalar(8.0); - model_builder->EndNode(); - model_builder->EndTree(); - - // TREE 3 - model_builder->StartTree(); - // Node P - model_builder->StartNode(0); - model_builder->LeafScalar(9.0); - model_builder->EndNode(); - model_builder->EndTree(); - - // TREE 4 - model_builder->StartTree(); - // Node Q - model_builder->StartNode(0); - model_builder->CategoricalTest(5, true, std::vector{0, 7}, false, 2, 1); - model_builder->EndNode(); - // Node R - model_builder->StartNode(1); - model_builder->LeafScalar(10.0); - model_builder->EndNode(); - // Node S - model_builder->StartNode(2); - model_builder->LeafScalar(11.0); - model_builder->EndNode(); - model_builder->EndTree(); - - // TREE 5 - model_builder->StartTree(); - // Node T - model_builder->StartNode(0); - model_builder->CategoricalTest(6, true, std::vector{1, 3}, true, 1, 2); - model_builder->EndNode(); - // Node U - model_builder->StartNode(1); - model_builder->LeafScalar(12.0); - model_builder->EndNode(); - // Node V - model_builder->StartNode(2); - model_builder->CategoricalTest(6, true, std::vector{1, 4}, true, 3, 4); - model_builder->EndNode(); - // Node W - model_builder->StartNode(3); - model_builder->CategoricalTest(6, true, std::vector{3}, true, 5, 6); - model_builder->EndNode(); - // Node X - model_builder->StartNode(4); - model_builder->LeafScalar(13.0); - model_builder->EndNode(); - // Node Y - model_builder->StartNode(5); - model_builder->LeafScalar(14.0); - model_builder->EndNode(); - // Node Z - model_builder->StartNode(6); - model_builder->LeafScalar(15.0); - model_builder->EndNode(); - model_builder->EndTree(); - return model_builder->CommitModel(); -}(); - -struct treelite_traversal_results { - std::vector feature_or_output; - std::vector depth; - std::vector parents; - std::vector tree_indices; -}; - -auto static const TRAVERSAL_RESULTS = - std::vector>{ - // Order: ABDEFGCHIJKLMNOPQRSTUVWYZX - std::make_pair(forest_order::depth_first, - treelite_traversal_results{ - // Order: {A, B, D, E, F, G, C, H, I, J, K, L, M, N, O, P, Q, R, S, - // T, U, V, W, Y, Z, X} - std::vector{0, 1, 1, 3, 2, 3, 0, 0, 4, 1, 5, 6, 5, - 7, 8, 9, 5, 10, 11, 6, 12, 6, 6, 14, 15, 13}, - // Order: {A, B, D, E, F, G, C, H, I, J, K, L, M, N, O, P, Q, R, - // S, T, U, V, W, Y, Z, X} - std::vector{0, 1, 2, 2, 3, 3, 1, 0, 1, 1, 2, 2, 0, - 1, 1, 0, 0, 1, 1, 0, 1, 1, 2, 3, 3, 2}, - // Order: {A, B, D, E, F, G, C, H, I, J, K, L, M, N, O, P, Q, - // R, S, T, U, V, W, Y, Z, X} - std::vector{0, 0, 1, 1, 3, 3, 0, 7, 7, 7, 9, 9, 12, - 12, 12, 15, 16, 16, 16, 19, 19, 19, 21, 22, 22, 21}, - // Order: {A, B, D, E, F, G, C, H, I, J, K, L, M, N, O, P, Q, R, - // S, T, U, V, W, Y, Z, X} - std::vector{0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 2, - 2, 2, 3, 4, 4, 4, 5, 5, 5, 5, 5, 5, 5}}), - // Order: ABCDEFGHIJKLMNOPQRSTUVWXYZ - std::make_pair(forest_order::breadth_first, - treelite_traversal_results{ - std::vector{0, 1, 0, 1, 3, 2, 3, 0, 4, 1, 5, 6, 5, - 7, 8, 9, 5, 10, 11, 6, 12, 6, 6, 13, 14, 15}, - std::vector{0, 1, 1, 2, 2, 3, 3, 0, 1, 1, 2, 2, 0, - 1, 1, 0, 0, 1, 1, 0, 1, 1, 2, 2, 3, 3}, - std::vector{0, 0, 0, 1, 1, 4, 4, 7, 7, 7, 9, 9, 12, - 12, 12, 15, 16, 16, 16, 19, 19, 19, 21, 21, 22, 22}, - std::vector{0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 2, - 2, 2, 3, 4, 4, 4, 5, 5, 5, 5, 5, 5, 5}}), - // Order: AHMPQTBCIJNORSUVDEKLWXFGYZ - std::make_pair(forest_order::layered_children_together, - treelite_traversal_results{ - // Order: {A, H, M, P, Q, T, B, C, I, J, N, O, R, S, U, V, D, E, K, - // L, W, X, F, G, Y, Z} - std::vector{0, 0, 5, 9, 5, 6, 1, 0, 4, 1, 7, 8, 10, - 11, 12, 6, 1, 3, 5, 6, 6, 13, 2, 3, 14, 15}, - // Order: {A, H, M, P, Q, T, B, C, I, J, N, O, R, S, U, V, D, E, - // K, L, W, X, F, G, Y, Z} - std::vector{0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 2, 2, 2, 2, 2, 2, 3, 3, 3, 3}, - // Order: {A, H, M, P, Q, T, B, C, I, J, N, O, R, S, U, V, D, E, - // K, L, W, X, F, G, Y, Z} - std::vector{0, 1, 2, 3, 4, 5, 0, 0, 1, 1, 2, 2, 4, - 4, 5, 5, 6, 6, 9, 9, 15, 15, 17, 17, 20, 20}, - // Order: {A, H, M, P, Q, T, B, C, I, J, N, O, R, S, U, V, D, E, - // K, L, W, X, F, G, Y, Z} - std::vector{0, 1, 2, 3, 4, 5, 0, 0, 1, 1, 2, 2, 4, - 4, 5, 5, 0, 0, 1, 1, 5, 5, 0, 0, 5, 5}}), - // Order: AHMPQTBINRUCJOSVDKWELXYFZG - std::make_pair(forest_order::layered_children_segregated, - treelite_traversal_results{ - // Order: {A, H, M, P, Q, T, B, I, N, R, U, C, J, O, S, V, D, K, W, - // E, L, X, Y, F, Z, G} - std::vector{0, 0, 5, 9, 5, 6, 1, 4, 7, 10, 12, 0, 1, - 8, 11, 6, 1, 5, 6, 3, 6, 13, 14, 2, 15, 3}, - // Order: {A, H, M, P, Q, T, B, I, N, R, U, C, J, O, S, V, D, K, - // W, E, L, X, Y, F, Z, G} - std::vector{0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 2, 2, 2, 2, 2, 2, 3, 3, 3, 3}, - // Order: {A, H, M, P, Q, T, B, I, N, R, U, C, J, O, S, V, D, K, - // W, E, L, X, Y, F, Z, G} - std::vector{0, 1, 2, 3, 4, 5, 0, 1, 2, 4, 5, 0, 1, - 2, 4, 5, 6, 12, 15, 6, 12, 15, 18, 19, 18, 19}, - // Order: {A, H, M, P, Q, T, B, I, N, R, U, C, J, O, S, V, D, K, - // W, E, L, X, Y, F, Z, G} - std::vector{0, 1, 2, 3, 4, 5, 0, 1, 2, 4, 5, 0, 1, - 2, 4, 5, 0, 1, 5, 0, 1, 5, 5, 0, 5, 0}}), - }; - -template -auto get_expected_for_each_result() -{ - return std::find_if(std::begin(TRAVERSAL_RESULTS), - std::end(TRAVERSAL_RESULTS), - [](auto&& pair) { return pair.first == order; }) - ->second; -} - -template -auto get_feature_or_outputs() -{ - auto result = std::vector{}; - node_transform(*SAMPLE_FOREST, - std::back_inserter(result), - [](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - auto feature_or_output = double{}; - if (node.is_leaf()) { - feature_or_output = node.get_output()[0]; - } else { - feature_or_output = node.get_feature(); - } - return feature_or_output; - }); - return result; -} - -template -auto get_depths() -{ - auto result = std::vector{}; - node_transform( - *SAMPLE_FOREST, - std::back_inserter(result), - [](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { return depth; }); - return result; -} - -template -auto get_parents() -{ - auto result = std::vector{}; - node_transform( - *SAMPLE_FOREST, - std::back_inserter(result), - [](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { return parent_index; }); - return result; -} - -template -auto get_tree_indices() -{ - auto result = std::vector{}; - node_transform( - *SAMPLE_FOREST, - std::back_inserter(result), - [](auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { return tree_id; }); - return result; -} - -template -auto get_categorical_count() -{ - return node_accumulate( - *SAMPLE_FOREST, - std::size_t{}, - [](auto&& acc, auto&& tree_id, auto&& node, auto&& depth, auto&& parent_index) { - return acc + node.is_categorical(); - }); -} - -TEST(ForestTraversal, depth_first) -{ - auto feature_or_output = get_feature_or_outputs(); - auto depths = get_depths(); - auto parents = get_parents(); - auto tree_indices = get_tree_indices(); - auto expected = get_expected_for_each_result(); - EXPECT_EQ(get_categorical_count(), SAMPLE_CATEGORICAL_COUNT); - for (auto i = std::size_t{}; i < expected.feature_or_output.size(); ++i) { - EXPECT_EQ(feature_or_output[i], expected.feature_or_output[i]); - } - for (auto i = std::size_t{}; i < expected.depth.size(); ++i) { - EXPECT_EQ(depths[i], expected.depth[i]); - } - for (auto i = std::size_t{}; i < expected.parents.size(); ++i) { - EXPECT_EQ(parents[i], expected.parents[i]); - } - for (auto i = std::size_t{}; i < expected.tree_indices.size(); ++i) { - EXPECT_EQ(tree_indices[i], expected.tree_indices[i]); - } -} - -TEST(ForestTraversal, breadth_first) -{ - auto feature_or_output = get_feature_or_outputs(); - auto depths = get_depths(); - auto parents = get_parents(); - auto tree_indices = get_tree_indices(); - auto expected = get_expected_for_each_result(); - EXPECT_EQ(get_categorical_count(), SAMPLE_CATEGORICAL_COUNT); - for (auto i = std::size_t{}; i < expected.feature_or_output.size(); ++i) { - EXPECT_EQ(feature_or_output[i], expected.feature_or_output[i]); - } - for (auto i = std::size_t{}; i < expected.depth.size(); ++i) { - EXPECT_EQ(depths[i], expected.depth[i]); - } - for (auto i = std::size_t{}; i < expected.parents.size(); ++i) { - EXPECT_EQ(parents[i], expected.parents[i]); - } - for (auto i = std::size_t{}; i < expected.tree_indices.size(); ++i) { - EXPECT_EQ(tree_indices[i], expected.tree_indices[i]); - } -} - -TEST(ForestTraversal, layered_children_segregated) -{ - auto feature_or_output = get_feature_or_outputs(); - auto depths = get_depths(); - auto parents = get_parents(); - auto tree_indices = get_tree_indices(); - auto expected = get_expected_for_each_result(); - EXPECT_EQ(get_categorical_count(), - SAMPLE_CATEGORICAL_COUNT); - for (auto i = std::size_t{}; i < expected.feature_or_output.size(); ++i) { - EXPECT_EQ(feature_or_output[i], expected.feature_or_output[i]); - } - for (auto i = std::size_t{}; i < expected.depth.size(); ++i) { - EXPECT_EQ(depths[i], expected.depth[i]); - } - for (auto i = std::size_t{}; i < expected.parents.size(); ++i) { - EXPECT_EQ(parents[i], expected.parents[i]); - } - for (auto i = std::size_t{}; i < expected.tree_indices.size(); ++i) { - EXPECT_EQ(tree_indices[i], expected.tree_indices[i]); - } -} - -TEST(ForestTraversal, layered_children_together) -{ - auto feature_or_output = get_feature_or_outputs(); - auto depths = get_depths(); - auto parents = get_parents(); - auto tree_indices = get_tree_indices(); - auto expected = get_expected_for_each_result(); - EXPECT_EQ(get_categorical_count(), - SAMPLE_CATEGORICAL_COUNT); - for (auto i = std::size_t{}; i < expected.feature_or_output.size(); ++i) { - EXPECT_EQ(feature_or_output[i], expected.feature_or_output[i]); - } - for (auto i = std::size_t{}; i < expected.depth.size(); ++i) { - EXPECT_EQ(depths[i], expected.depth[i]); - } - for (auto i = std::size_t{}; i < expected.parents.size(); ++i) { - EXPECT_EQ(parents[i], expected.parents[i]); - } - for (auto i = std::size_t{}; i < expected.tree_indices.size(); ++i) { - EXPECT_EQ(tree_indices[i], expected.tree_indices[i]); - } -} - -} // namespace forest -} // namespace ML diff --git a/cpp/tests/sg/rf_test.cu b/cpp/tests/sg/rf_test.cu index 9ae8c932f1..8e769d2f14 100644 --- a/cpp/tests/sg/rf_test.cu +++ b/cpp/tests/sg/rf_test.cu @@ -5,10 +5,6 @@ #include #include #include -#include -#include -#include -#include #include #include @@ -37,6 +33,10 @@ #include #include #include +#include +#include +#include +#include #include #include @@ -177,7 +177,7 @@ std::ostream& operator<<(std::ostream& os, const RfTestParams& ps) } template -std::shared_ptr> FilPredict( +std::shared_ptr> nvForestPredict( const raft::handle_t& handle, RfTestParams params, DataT* X_transpose, @@ -186,12 +186,12 @@ std::shared_ptr> FilPredict( auto pred = std::shared_ptr>(); auto workspace = std::shared_ptr>(); // Scratch space if constexpr (std::is_integral_v) { - // For classifiers, allocate extra scratch space to store probabilities from FIL + // For classifiers, allocate extra scratch space to store probabilities from nvForest // We will perform argmax to convert probabilities into class outputs. pred = std::make_shared>(params.n_rows); workspace = std::make_shared>(params.n_rows * params.n_labels); } else { - // For regressors, no need to post-process predictions from FIL + // For regressors, no need to post-process predictions from nvForest static_assert(std::is_same_v, "LabelT and DataT must be identical for regression task"); pred = std::make_shared>(params.n_rows); @@ -200,25 +200,25 @@ std::shared_ptr> FilPredict( TreeliteModelHandle model; build_treelite_forest(&model, forest, params.n_cols); - auto fil_model = ML::fil::import_from_treelite_handle(model, - ML::fil::tree_layout::breadth_first, - 128, - std::is_same_v, - raft_proto::device_type::gpu, - handle.get_device(), - handle.get_next_usable_stream()); + auto nvforest_model = nvforest::import_from_treelite_handle(model, + nvforest::tree_layout::breadth_first, + 128, + std::is_same_v, + raft_proto::device_type::gpu, + handle.get_device(), + handle.get_next_usable_stream()); handle.sync_stream(); handle.sync_stream_pool(); delete static_cast(model); - fil_model.predict(handle, - workspace->data().get(), - X_transpose, - params.n_rows, - raft_proto::device_type::gpu, - raft_proto::device_type::gpu, - ML::fil::infer_kind::default_kind, - 1); + nvforest_model.predict(handle, + workspace->data().get(), + X_transpose, + params.n_rows, + raft_proto::device_type::gpu, + raft_proto::device_type::gpu, + nvforest::infer_kind::default_kind, + 1); handle.sync_stream(); handle.sync_stream_pool(); @@ -259,10 +259,10 @@ std::shared_ptr> FilPredict( } template -auto FilPredictProba(const raft::handle_t& handle, - RfTestParams params, - DataT* X_transpose, - RandomForestMetaData* forest) +auto nvForestPredictProba(const raft::handle_t& handle, + RfTestParams params, + DataT* X_transpose, + RandomForestMetaData* forest) { static_assert(std::is_integral_v, "Must be classification"); @@ -271,25 +271,25 @@ auto FilPredictProba(const raft::handle_t& handle, TreeliteModelHandle model; build_treelite_forest(&model, forest, params.n_cols); - auto fil_model = ML::fil::import_from_treelite_handle(model, - ML::fil::tree_layout::breadth_first, - 128, - std::is_same_v, - raft_proto::device_type::gpu, - handle.get_device(), - handle.get_next_usable_stream()); + auto nvforest_model = nvforest::import_from_treelite_handle(model, + nvforest::tree_layout::breadth_first, + 128, + std::is_same_v, + raft_proto::device_type::gpu, + handle.get_device(), + handle.get_next_usable_stream()); handle.sync_stream(); handle.sync_stream_pool(); delete static_cast(model); - fil_model.predict(handle, - pred->data().get(), - X_transpose, - params.n_rows, - raft_proto::device_type::gpu, - raft_proto::device_type::gpu, - ML::fil::infer_kind::default_kind, - 1); + nvforest_model.predict(handle, + pred->data().get(), + X_transpose, + params.n_rows, + raft_proto::device_type::gpu, + raft_proto::device_type::gpu, + nvforest::infer_kind::default_kind, + 1); handle.sync_stream(); handle.sync_stream_pool(); @@ -495,37 +495,38 @@ class RfSpecialisedTest { return std::abs(max_element - second_max_element); } - // Compare fil against native rf predictions + // Compare nvForest against native rf predictions // Only for single precision models - void TestFilPredict() + void TestNvForestPredict() { if constexpr (std::is_same_v) { return; } else { auto stream_pool = std::make_shared(params.n_streams); raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); - auto fil_pred = FilPredict(handle, params, X_transpose.data().get(), forest.get()); + auto nvforest_pred = nvForestPredict(handle, params, X_transpose.data().get(), forest.get()); - thrust::host_vector h_fil_pred(*fil_pred); + thrust::host_vector h_nvforest_pred(*nvforest_pred); thrust::host_vector h_pred(*predictions); - thrust::host_vector h_fil_pred_prob; + thrust::host_vector h_nvforest_pred_prob; if constexpr (std::is_integral_v) { - h_fil_pred_prob = *FilPredictProba(handle, params, X_transpose.data().get(), forest.get()); + h_nvforest_pred_prob = + *nvForestPredictProba(handle, params, X_transpose.data().get(), forest.get()); } float tol = 1e-2; - for (std::size_t i = 0; i < h_fil_pred.size(); i++) { + for (std::size_t i = 0; i < h_nvforest_pred.size(); i++) { // If the output probabilities are very similar for different classes - // FIL may output a different class due to numerical differences + // nvForest may output a different class due to numerical differences // Skip these cases if constexpr (std::is_integral_v) { int num_outputs = forest->trees[0]->num_outputs; - auto min_diff = MinDifference(&h_fil_pred_prob[i * num_outputs], num_outputs); + auto min_diff = MinDifference(&h_nvforest_pred_prob[i * num_outputs], num_outputs); if (min_diff < tol) continue; } - EXPECT_LE(abs(h_fil_pred[i] - h_pred[i]), tol); + EXPECT_LE(abs(h_nvforest_pred[i] - h_pred[i]), tol); } } } @@ -570,7 +571,7 @@ class RfSpecialisedTest { TestMinImpurity(); TestTreeSize(); TestInstanceCounts(); - TestFilPredict(); + TestNvForestPredict(); TestFeatureImportances(); } @@ -677,30 +678,30 @@ TEST(RfTests, IntegerOverflow) // Check we have actually learned something EXPECT_GT(forest->trees[0]->leaf_counter, 1); - // See if FIL overflows + // See if nvForest overflows thrust::device_vector pred(m); TreeliteModelHandle model; build_treelite_forest(&model, forest_ptr, n); - auto fil_model = ML::fil::import_from_treelite_handle(model, - ML::fil::tree_layout::breadth_first, - 128, - false, - raft_proto::device_type::gpu, - handle.get_device(), - handle.get_next_usable_stream()); + auto nvforest_model = nvforest::import_from_treelite_handle(model, + nvforest::tree_layout::breadth_first, + 128, + false, + raft_proto::device_type::gpu, + handle.get_device(), + handle.get_next_usable_stream()); handle.sync_stream(); handle.sync_stream_pool(); delete static_cast(model); - fil_model.predict(handle, - pred.data().get(), - X.data().get(), - m, - raft_proto::device_type::gpu, - raft_proto::device_type::gpu, - ML::fil::infer_kind::default_kind, - 1); + nvforest_model.predict(handle, + pred.data().get(), + X.data().get(), + m, + raft_proto::device_type::gpu, + raft_proto::device_type::gpu, + nvforest::infer_kind::default_kind, + 1); handle.sync_stream(); handle.sync_stream_pool(); } diff --git a/dependencies.yaml b/dependencies.yaml index 12705a4401..3d3106ca79 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -19,9 +19,11 @@ files: - depends_on_dask_cuda - depends_on_dask_cudf - depends_on_libcuvs + - depends_on_libnvforest - depends_on_libraft - depends_on_librmm - depends_on_numba_cuda + - depends_on_nvforest - depends_on_pylibraft - depends_on_raft_dask - depends_on_rapids_logger @@ -49,9 +51,11 @@ files: - depends_on_dask_cuda - depends_on_dask_cudf - depends_on_libcuvs + - depends_on_libnvforest - depends_on_libraft - depends_on_librmm - depends_on_numba_cuda + - depends_on_nvforest - depends_on_pylibraft - depends_on_raft_dask - depends_on_rapids_logger @@ -73,6 +77,7 @@ files: - cuda - cuda_version - depends_on_libcuvs + - depends_on_libnvforest - depends_on_libraft_headers - depends_on_librmm checks: @@ -91,6 +96,7 @@ files: - cuda - cuda_version - depends_on_libcuvs + - depends_on_libnvforest - depends_on_libraft_headers - depends_on_librmm docs: @@ -101,6 +107,7 @@ files: - py_version - depends_on_cuml - depends_on_libcuml + - depends_on_libnvforest - depends_on_rapids_dask_dependency - depends_on_dask_cuda - depends_on_dask_cudf @@ -110,6 +117,7 @@ files: includes: - cuda_version - depends_on_libcuml + - depends_on_libnvforest - test_libcuml - test_cpp test_python: @@ -119,6 +127,7 @@ files: - depends_on_cuml - depends_on_cupy - depends_on_libcuml + - depends_on_nvforest - py_version - test_python - test_python_xgboost @@ -129,6 +138,7 @@ files: - depends_on_cuml - depends_on_cupy - depends_on_libcuml + - depends_on_libnvforest - py_version - test_python - test_python_xgboost @@ -147,6 +157,7 @@ files: - depends_on_dask_cuda - depends_on_dask_cudf - depends_on_numba_cuda + - depends_on_nvforest - depends_on_pylibraft - depends_on_raft_dask - depends_on_rmm @@ -172,8 +183,10 @@ files: - depends_on_cuda_python - depends_on_libcuml - depends_on_libcuvs + - depends_on_libnvforest - depends_on_libraft - depends_on_librmm + - depends_on_nvforest - depends_on_pylibraft - depends_on_rmm - py_build_cuml @@ -189,6 +202,7 @@ files: - depends_on_cupy - depends_on_libcuml - depends_on_numba_cuda + - depends_on_nvforest - depends_on_pylibraft - depends_on_rmm - py_run_cuml @@ -236,6 +250,7 @@ files: includes: - common_build - depends_on_libcuvs + - depends_on_libnvforest - depends_on_libraft - depends_on_librmm - depends_on_rapids_logger @@ -247,6 +262,7 @@ files: includes: - cuda_wheels - depends_on_libcuvs + - depends_on_libnvforest - depends_on_libraft - depends_on_librmm - depends_on_rapids_logger @@ -854,6 +870,29 @@ dependencies: - output_types: conda packages: - &libcuvs_unsuffixed libcuvs==26.6.*,>=0.0.0a0 + depends_on_libnvforest: + common: + - output_types: conda + packages: + - &libnvforest_unsuffixed libnvforest==26.6.*,>=0.0.0a0 + - output_types: requirements + packages: + # pip recognizes the index as a global option for the requirements.txt file + - --extra-index-url=https://pypi.anaconda.org/rapidsai-wheels-nightly/simple + specific: + - output_types: [requirements, pyproject] + matrices: + - matrix: + cuda: "12.*" + cuda_suffixed: "true" + packages: + - libnvforest-cu12==26.6.*,>=0.0.0a0 + - matrix: + cuda: "13.*" + cuda_suffixed: "true" + packages: + - libnvforest-cu13==26.6.*,>=0.0.0a0 + - {matrix: null, packages: [*libnvforest_unsuffixed]} depends_on_libraft: common: - output_types: conda @@ -927,6 +966,31 @@ dependencies: - matrix: packages: - *numba_cuda + depends_on_nvforest: + common: + - output_types: conda + packages: + - &nvforest_unsuffixed nvforest==26.6.*,>=0.0.0a0 + - output_types: requirements + packages: + # pip recognizes the index as a global option for the requirements.txt file + - --extra-index-url=https://pypi.anaconda.org/rapidsai-wheels-nightly/simple + specific: + - output_types: [requirements, pyproject] + matrices: + - matrix: + cuda: "12.*" + cuda_suffixed: "true" + packages: + - nvforest-cu12==26.6.*,>=0.0.0a0 + - matrix: + cuda: "13.*" + cuda_suffixed: "true" + packages: + - nvforest-cu13==26.6.*,>=0.0.0a0 + - matrix: + packages: + - *nvforest_unsuffixed depends_on_pylibraft: common: - output_types: conda diff --git a/docs/source/FIL.rst b/docs/source/FIL.rst index 680295c1ab..6be62652ba 100644 --- a/docs/source/FIL.rst +++ b/docs/source/FIL.rst @@ -1,238 +1,13 @@ FIL - RAPIDS Forest Inference Library ===================================== +.. note:: + + This module is deprecated. New code should use + `nvForest `_ instead. + The Forest Inference Library (FIL) is a component of cuML, providing a high-performance inference engine designed to accelerate tree-based machine learning models on both GPU and CPU. FIL delivers significant speedups over traditional CPU-based inference while maintaining compatibility with models trained in popular frameworks. - -**Key Benefits:** - -- FIL typically offers a speedup of 80x or more over scikit-learn native execution -- Support for XGBoost, Scikit-Learn, LightGBM, and Treelite-compatible models -- Seamless GPU/CPU execution switching -- Built-in auto-optimization for maximum performance -- Advanced inference APIs for granular tree analysis - -**Quick Start:** - -.. code-block:: python - - import xgboost as xgb - import numpy as np - from cuml.fil import ForestInference - - # Train your model as usual and save it - xgb_model = xgb.XGBClassifier() - xgb_model.fit(X_train, y_train) - xgb_model.save_model("xgb_model.ubj") - - # Load into FIL and auto-tune for your batch size - fil_model = ForestInference.load("xgb_model.ubj", is_classifier=True) - fil_model.optimize(batch_size=1024) - - # Now you can predict with FIL directly - predictions = fil_model.predict(X_test) - probabilities = fil_model.predict_proba(X_test) - -Performance Optimization -------------------------- -FIL includes built-in auto-optimization that automatically tunes performance hyperparameters for your specific model and batch size, eliminating the need for manual tuning in most cases: - -.. code-block:: python - - fil_model = ForestInference.load("model.ubj", is_classifier=True) - fil_model.optimize(batch_size=1_000_000) - - # Check which hyperparameters were selected - print(f"Layout: {fil_model.layout}") - print(f"Chunk size: {fil_model.default_chunk_size}") - - result = fil_model.predict(data) - -The optimization process tests different memory layouts and chunk sizes to find the optimal configuration for your specific use case. - -**Key Hyperparameters:** - -- ``layout``: Determines the order in which tree nodes are arranged in memory (depth_first, layered, breadth_first) -- ``default_chunk_size``: Controls the granularity of parallelization during inference -- ``align_bytes``: Cache line alignment for optimal memory access patterns - -**Manual Tuning:** -For advanced users, you can experiment with the ``align_bytes`` parameter. Its default value is typically close enough to optimal that it is not automatically searched during auto-optimization, but to squeeze the most performance possible out of FIL, try either 0 or 128 on GPU and 0 or 64 on CPU. - -Optional CPU Execution ----------------------- -While FIL offers the most benefit for large models and batch sizes by taking advantage of the speed and parallelism of NVIDIA GPUs, it can also be used to speed up inference on CPUs. This can be convenient for testing in environments without access to GPUs. It can also be useful for deployments which experience dramatic shifts in traffic. When the number of incoming inference requests is low, CPU execution can be used. When traffic spikes, the deployment can seamlessly scale up onto GPUs in order to handle the additional load as cheaply as possible without significantly increasing latency. - -You can use FIL in CPU mode with a context manager: - -.. code-block:: python - - from cuml.fil import ForestInference, set_fil_device_type - - with set_fil_device_type("cpu"): - fil_model = ForestInference.load("xgboost_model.ubj") - result = fil_model.predict(data) - -Advanced Prediction APIs -------------------------- -FIL includes advanced prediction methods that provide granular information about individual trees in the ensemble, enabling novel ensembling techniques and analysis: - -**Per-Tree Predictions** -The ``.predict_per_tree`` method returns the output of every single tree individually: - -.. code-block:: python - - per_tree = fil_model.predict_per_tree(X) - mean = per_tree.mean(axis=1) - lower = np.percentile(per_tree, 10, axis=1) - upper = np.percentile(per_tree, 90, axis=1) - -This enables advanced techniques like: - -- Weighted voting based on tree age, out-of-bag AUC, or data-drift scores -- Prediction intervals without bootstrapping -- Novel ensembling techniques with no retraining required - -**Leaf Node Analysis** -The ``.apply`` method returns the leaf node ID for every tree, enabling similarity analysis: - -.. code-block:: python - - leaf = fil_model.apply(X) - sim = (leaf[i] == leaf[j]).mean() # fraction of matching leaves - print(f"{sim:.0%} of trees agree on rows {i} & {j}") - -This opens forest models to novel uses beyond straightforward regression or classification, such as measuring data similarity and understanding model behavior. - -Use Cases ---------- -FIL is ideal for many scenarios: - -**High-Performance Applications:** - -- User-facing APIs where every millisecond counts -- High-volume batch jobs (ad-click scoring, IoT analytics) -- Real-time inference with sub-10ms latency requirements - -**Flexible Deployment:** - -- Hybrid deployments - same model file, choose CPU or GPU at runtime -- Prototype locally and deploy to GPU-accelerated production servers -- Scale down to CPU-only machines during light traffic, scale up with GPUs during peak loads - -**Cost Optimization:** - -- One GPU can replace CPUs with 50+ cores -- Significant cost reduction for high-throughput inference workloads -- Efficient resource utilization across different traffic patterns - -**Advanced Analytics:** - -- Novel ensembling techniques with per-tree analysis -- Data similarity measurement and model interpretability -- Prediction intervals and uncertainty quantification - -API Reference -============= - -See the :doc:`API reference ` for the API documentation. - -Migration Guide -=============== - -FIL Redesign in RAPIDS 25.04 ------------------------------ -FIL was completely redesigned in RAPIDS 25.04 with a new C++ implementation that provides significant performance improvements and new features: - -**Key Changes in 25.04:** - -- New C++ implementation for batched inference on GPU and CPU -- Built-in auto-optimization with ``.optimize()`` method -- Advanced inference APIs (``.predict_per_tree``, ``.apply``) -- Up to 4x faster GPU throughput than previous versions -- Enhanced memory layouts and cache optimization -- New parameter structure (``layout``, ``align_bytes``) -- Moved ``threshold`` from ``.load()`` to ``.predict()`` - -Migration from RAPIDS 25.04 to 25.06 (Output Shape Changes) ------------------------------------------------------------ -In RAPIDS 25.06, the shape of output arrays changed for some models. Binary classifiers now return an array of solely the probabilities of the positive class for ``predict_proba`` calls. This both reduces memory requirements and improves performance. To convert to the old format, the following snippet can be used: - -.. code-block:: python - - import numpy as np # Use cupy or numpy depending on which you use for input data - - out = fil_model.predict_proba(input_data) - # Starting in RAPIDS 25.06, the following can be used to obtain the old output shape - out = np.stack([1 - out, out], axis=1) - -Additionally, ``.predict`` calls now output two-dimensional arrays beginning in 25.06. This is in preparation for supporting multi-target regression and classification models. The old shape can be obtained via the following snippet: - -.. code-block:: python - - import numpy as np # Use cupy or numpy depending on which you use for input data - - out = fil_model.predict(input_data) - # Starting in RAPIDS 25.06, the following can be used to obtain the old output shape - out = out.flatten() - -To use these new behaviors immediately, the ``ForestInference`` estimator can be imported from the ``experimental`` namespace: - -.. code-block:: python - - from cuml.experimental.fil import ForestInference - -Migration from RAPIDS 24.12 to 25.04 ------------------------------------- - -**Before (RAPIDS 24.12):** - -.. code-block:: python - - fil_model = ForestInference.load( - "./model.ubj", - is_classifier=True, - algo='TREE_REORG', # Deprecated - threshold=0.5, # Now moved to predict() - storage_type='DENSE' # Deprecated - ) - predictions = fil_model.predict(data) - -**After (RAPIDS 25.04):** - -.. code-block:: python - - fil_model = ForestInference.load( - "./model.ubj", - is_classifier=True, - layout='depth_first' # New parameter - ) - predictions = fil_model.predict(data, threshold=0.5) # threshold moved here - -Deprecated ``load`` Parameters -^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ -As of RAPIDS 25.04, the following hyperparameters accepted by the ``.load`` method of previous versions of FIL have been deprecated. - -- ``threshold`` (will trigger a deprecation warning if used; pass to ``.predict`` instead) -- ``algo`` (ignored, but a warning will be logged) -- ``storage_type`` (ignored, but a warning will be logged) -- ``blocks_per_sm`` (ignored, but a warning will be logged) -- ``threads_per_tree`` (ignored, but a warning will be logged) -- ``n_items`` (ignored, but a warning will be logged) -- ``compute_shape_str`` (ignored, but a warning will be logged) - -New ``load`` Parameters -^^^^^^^^^^^^^^^^^^^^^^^ -As of RAPIDS 25.04, the following new hyperparameters can be passed to the ``.load`` method - -- ``layout``: Replaces the functionality of ``algo`` and specifies the in-memory layout of nodes in FIL forests. One of ``'depth_first'`` (default), ``'layered'`` or ``'breadth_first'``. -- ``align_bytes``: If specified, trees will be padded such that their in-memory size is a multiple of this value. This can sometimes improve performance by guaranteeing that memory reads from trees begin on a cache line boundary. - -New Prediction Parameters -^^^^^^^^^^^^^^^^^^^^^^^^^ -As of RAPIDS 25.04, all prediction methods accept a ``chunk_size`` parameter, which determines how batches are further subdivided for parallel processing. The optimal value depends on hardware, model, and batch size, and it is difficult to predict in advance. Typically, it is best to use the ``.optimize`` method to determine the best chunk size for a given batch size. If ``chunk_size`` must be set manually, the only general rule of thumb is that larger batch sizes generally benefit from larger chunk sizes. On GPU, ``chunk_size`` can be any power of 2 from 1 to 32. On CPU, ``chunk_size`` can be any power of 2, but values above 512 rarely offer any benefit. - -Additionally, ``threshold`` has been converted from a ``.load`` parameter to a ``.predict`` parameter. diff --git a/python/cuml/CMakeLists.txt b/python/cuml/CMakeLists.txt index 3a186478b2..97951b7d93 100644 --- a/python/cuml/CMakeLists.txt +++ b/python/cuml/CMakeLists.txt @@ -95,7 +95,6 @@ add_subdirectory(cuml/datasets) add_subdirectory(cuml/decomposition) add_subdirectory(cuml/ensemble) add_subdirectory(cuml/explainer) -add_subdirectory(cuml/fil) add_subdirectory(cuml/linear_model) add_subdirectory(cuml/manifold) add_subdirectory(cuml/metrics) diff --git a/python/cuml/cuml/__init__.py b/python/cuml/cuml/__init__.py index f60e531e23..6a5c71485e 100644 --- a/python/cuml/cuml/__init__.py +++ b/python/cuml/cuml/__init__.py @@ -36,7 +36,7 @@ from cuml.explainer.kernel_shap import KernelExplainer from cuml.explainer.permutation_shap import PermutationExplainer from cuml.explainer.tree_shap import TreeExplainer -from cuml.fil import ForestInference, fil +from cuml.fil import ForestInference from cuml.internals.base import Base from cuml.internals.global_settings import ( GlobalSettings, @@ -95,7 +95,6 @@ def __getattr__(name): # Modules "common", "feature_extraction", - "fil", "metrics", "multiclass", "naive_bayes", diff --git a/python/cuml/cuml/benchmark/algorithms.py b/python/cuml/cuml/benchmark/algorithms.py index 5b4784ec83..03b1f09476 100644 --- a/python/cuml/cuml/benchmark/algorithms.py +++ b/python/cuml/cuml/benchmark/algorithms.py @@ -32,7 +32,6 @@ fit_kneighbors, fit_predict, fit_transform, - predict, transform, ) from cuml.benchmark.gpu_check import is_cuml_available @@ -45,7 +44,6 @@ fit_kneighbors, fit_predict, fit_transform, - predict, transform, ) from gpu_check import is_cuml_available # noqa: E402 @@ -57,12 +55,7 @@ # GPU-specific helper functions (only available when GPU libs present) _build_cpu_skl_classifier = None -_build_fil_classifier = None -_build_fil_skl_classifier = None -_build_gtil_classifier = None _build_mnmg_umap = None -_build_optimized_fil_classifier = None -_treelite_fil_accuracy_score = None # cuML preprocessing classes (fallback to sklearn if not available) MaxAbsScaler = sklearn.preprocessing.MaxAbsScaler @@ -91,16 +84,11 @@ cuml = _cuml cuml_metrics = _cuml_metrics - # Import GPU-specific helper functions (package path; standalone uses except above) - from cuml.benchmark.bench_helper_funcs import ( - _build_cpu_skl_classifier, - _build_fil_classifier, - _build_fil_skl_classifier, - _build_gtil_classifier, - _build_mnmg_umap, - _build_optimized_fil_classifier, - _treelite_fil_accuracy_score, - ) + # Import GPU-specific helper functions (support package + standalone execution) + try: + from cuml.benchmark.bench_helper_funcs import _build_mnmg_umap + except ImportError: + from bench_helper_funcs import _build_mnmg_umap # noqa: E402 # Optional treelite import try: @@ -346,7 +334,6 @@ def all_algorithms(): cuml_KNeighborsRegressor = cuml.neighbors.KNeighborsRegressor cuml_MultinomialNB = cuml.naive_bayes.MultinomialNB cuml_UMAP = cuml.manifold.UMAP - cuml_ForestInference = cuml.ForestInference accuracy_fn = cuml_metrics.accuracy_score r2_fn = cuml_metrics.r2_score trustworthiness_fn = cuml_metrics.trustworthiness @@ -363,7 +350,7 @@ def all_algorithms(): None ) cuml_KNeighborsClassifier = cuml_KNeighborsRegressor = None - cuml_MultinomialNB = cuml_UMAP = cuml_ForestInference = None + cuml_MultinomialNB = cuml_UMAP = None accuracy_fn = metrics.accuracy_score r2_fn = metrics.r2_score trustworthiness_fn = None @@ -713,70 +700,6 @@ def all_algorithms(): ) ) - # Add FIL algorithms if treelite and cuML are available - if ( - is_cuml_available() - and treelite is not None - and _build_fil_classifier is not None - ): - algorithms.extend( - [ - AlgorithmPair( - treelite, - cuml_ForestInference, - shared_args=dict(num_rounds=100, max_depth=10), - cuml_args=dict( - is_classifier=False, - threshold=0.5, - precision="float32", - layout="depth_first", - ), - name="FIL", - accepts_labels=False, - setup_cpu_func=_build_gtil_classifier, - setup_cuml_func=_build_fil_classifier, - cpu_data_prep_hook=_treelite_format_hook, - accuracy_function=_treelite_fil_accuracy_score, - bench_func=predict, - ), - AlgorithmPair( - treelite, - cuml_ForestInference, - shared_args=dict(n_estimators=100, max_leaf_nodes=2**10), - cuml_args=dict( - is_classifier=False, - threshold=0.5, - precision="float32", - layout="depth_first", - ), - name="Sparse-FIL-SKL", - accepts_labels=False, - setup_cpu_func=_build_cpu_skl_classifier, - setup_cuml_func=_build_fil_skl_classifier, - accuracy_function=_treelite_fil_accuracy_score, - bench_func=predict, - ), - AlgorithmPair( - treelite, - cuml_ForestInference, - shared_args=dict(num_rounds=100, max_depth=10), - cuml_args=dict( - is_classifier=False, - threshold=0.5, - precision="float32", - layout="depth_first", - ), - name="FIL-Optimized", - accepts_labels=False, - setup_cpu_func=_build_gtil_classifier, - setup_cuml_func=_build_optimized_fil_classifier, - cpu_data_prep_hook=_treelite_format_hook, - accuracy_function=_treelite_fil_accuracy_score, - bench_func=predict, - ), - ] - ) - # Add MNMG (multi-node multi-GPU) algorithms if cuML.dask is available if is_cuml_available(): try: diff --git a/python/cuml/cuml/benchmark/bench_helper_funcs.py b/python/cuml/cuml/benchmark/bench_helper_funcs.py index cdf7ee5202..5d54128b56 100644 --- a/python/cuml/cuml/benchmark/bench_helper_funcs.py +++ b/python/cuml/cuml/benchmark/bench_helper_funcs.py @@ -2,15 +2,11 @@ # SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # -import os import pickle as pickle import sys -from time import perf_counter import numpy as np import pandas as pd -import sklearn.ensemble as skl_ensemble -from sklearn import metrics as sklearn_metrics # Supports both package and standalone execution try: @@ -28,8 +24,6 @@ cuda = None cuml = None input_utils = None -get_fil_device_type = None -set_fil_device_type = None DeviceType = None UMAP = None @@ -38,10 +32,6 @@ import cupy as cp from numba import cuda - import cuml - from cuml.fil import get_fil_device_type, set_fil_device_type - from cuml.internals import input_utils - from cuml.internals.device_type import DeviceType from cuml.manifold import UMAP @@ -128,291 +118,6 @@ def _training_data_to_numpy(X, y): return X_np, y_np -def _build_fil_classifier(m, data, args, tmpdir): - """Setup function for FIL classification benchmarking. - - Note: This function requires GPU libraries (cuML) to be available. - """ - if not is_cuml_available(): - raise RuntimeError( - "FIL classifier requires GPU libraries (cuML). " - "Not available in CPU-only mode." - ) - import xgboost as xgb - - train_data, train_label = _training_data_to_numpy(data[0], data[1]) - - dtrain = xgb.DMatrix(train_data, label=train_label) - - params = { - "silent": 1, - "eval_metric": "error", - "objective": "binary:logistic", - "tree_method": "hist", - "device": "cuda", - } - params.update(args) - max_depth = args["max_depth"] - num_rounds = args["num_rounds"] - n_feature = data[0].shape[1] - train_size = data[0].shape[0] - model_name = f"xgb_{max_depth}_{num_rounds}_{n_feature}_{train_size}.ubj" - model_path = os.path.join(tmpdir, model_name) - bst = xgb.train(params, dtrain, num_rounds) - bst.save_model(model_path) - - fil_kwargs = { - param: args[input_name] - for param, input_name in ( - ("is_classifier", "is_classifier"), - ("threshold", "threshold"), - ("precision", "precision"), - ("layout", "layout"), - ) - if input_name in args - } - - return m.load(model_path, **fil_kwargs) - - -class OptimizedFilWrapper: - """Helper class to make use of optimized parameters in FIL""" - - def __init__(self, fil_model, optimal_chunk_size, infer_type="default"): - self.fil_model = fil_model - self.predict_kwargs = {"chunk_size": optimal_chunk_size} - self.infer_type = infer_type - - def predict(self, X): - if self.infer_type == "per_tree": - return self.fil_model.predict_per_tree(X, **self.predict_kwargs) - return self.fil_model.predict(X, **self.predict_kwargs) - - -def _build_optimized_fil_classifier(m, data, args, tmpdir): - """Setup function for FIL classification benchmarking with optimal - parameters. - - Note: This function requires GPU libraries (cuML) to be available. - """ - if not is_cuml_available(): - raise RuntimeError( - "Optimized FIL classifier requires GPU libraries (cuML). " - "Not available in CPU-only mode." - ) - import xgboost as xgb - - with set_fil_device_type("gpu"): - train_data, train_label = _training_data_to_numpy(data[0], data[1]) - - dtrain = xgb.DMatrix(train_data, label=train_label) - - params = { - "silent": 1, - "eval_metric": "error", - "objective": "binary:logistic", - "tree_method": "hist", - "device": "cuda", - } - params.update(args) - max_depth = args["max_depth"] - num_rounds = args["num_rounds"] - n_feature = data[0].shape[1] - train_size = data[0].shape[0] - model_name = ( - f"xgb_{max_depth}_{num_rounds}_{n_feature}_{train_size}.ubj" - ) - model_path = os.path.join(tmpdir, model_name) - bst = xgb.train(params, dtrain, num_rounds) - bst.save_model(model_path) - - allowed_chunk_sizes = [1, 2, 4, 8, 16, 32] - if get_fil_device_type() is DeviceType.host: - allowed_chunk_sizes.extend((64, 128, 256)) - - fil_kwargs = { - param: args[input_name] - for param, input_name in ( - ("is_classifier", "is_classifier"), - ("threshold", "threshold"), - ("precision", "precision"), - ("layout", "layout"), - ) - if input_name in args - } - infer_type = args.get("infer_type", "default") - - optimal_layout = "breadth_first" - optimal_chunk_size = 1 - best_time = None - optimization_cycles = 5 - - allowed_layout_types = ["breadth_first", "depth_first", "layered"] - for layout in allowed_layout_types: - fil_kwargs["layout"] = layout - for chunk_size in allowed_chunk_sizes: - call_args = {"chunk_size": chunk_size} - fil_model = m.load(model_path, **fil_kwargs) - if infer_type == "per_tree": - fil_model.predict_per_tree(train_data, **call_args) - else: - fil_model.predict(train_data, **call_args) - begin = perf_counter() - if infer_type == "per_tree": - fil_model.predict_per_tree(train_data, **call_args) - else: - for _ in range(optimization_cycles): - fil_model.predict(train_data, **call_args) - end = perf_counter() - elapsed = end - begin - if best_time is None or elapsed < best_time: - best_time = elapsed - optimal_chunk_size = chunk_size - optimal_layout = layout - - fil_kwargs["layout"] = optimal_layout - - return OptimizedFilWrapper( - m.load(model_path, **fil_kwargs), - optimal_chunk_size, - infer_type=infer_type, - ) - - -def _build_fil_skl_classifier(m, data, args, tmpdir): - """Trains an SKLearn classifier and returns a FIL version of it. - - Note: This function requires GPU libraries (cuML) to be available. - """ - if not is_cuml_available(): - raise RuntimeError( - "FIL SKLearn classifier requires GPU libraries (cuML). " - "Not available in CPU-only mode." - ) - - train_data, train_label = _training_data_to_numpy(data[0], data[1]) - - params = { - "n_estimators": 100, - "max_leaf_nodes": 2**10, - "max_features": "sqrt", - "n_jobs": -1, - "random_state": 42, - } - params.update(args) - - # remove keyword arguments not understood by SKLearn - for param_name in [ - "is_classifier", - "threshold", - "precision", - "layout", - ]: - params.pop(param_name, None) - - max_leaf_nodes = args["max_leaf_nodes"] - n_estimators = args["n_estimators"] - n_feature = data[0].shape[1] - train_size = data[0].shape[0] - model_name = ( - f"skl_{max_leaf_nodes}_{n_estimators}_{n_feature}_" - + f"{train_size}.model.pkl" - ) - model_path = os.path.join(tmpdir, model_name) - skl_model = skl_ensemble.RandomForestClassifier(**params) - skl_model.fit(train_data, train_label) - pickle.dump(skl_model, open(model_path, "wb")) - - fil_kwargs = { - param: args[input_name] - for param, input_name in ( - ("is_classifier", "is_classifier"), - ("threshold", "threshold"), - ("precision", "precision"), - ("layout", "layout"), - ) - if input_name in args - } - - return m.load_from_sklearn(skl_model, **fil_kwargs) - - -def _build_cpu_skl_classifier(m, data, args, tmpdir): - """Loads the SKLearn classifier and returns it""" - - max_leaf_nodes = args["max_leaf_nodes"] - n_estimators = args["n_estimators"] - n_feature = data[0].shape[1] - train_size = data[0].shape[0] - model_name = ( - f"skl_{max_leaf_nodes}_{n_estimators}_{n_feature}_" - + f"{train_size}.model.pkl" - ) - model_path = os.path.join(tmpdir, model_name) - - skl_model = pickle.load(open(model_path, "rb")) - return skl_model - - -class GtilWrapper: - """Helper class to provide interface to GTIL compatible with - benchmarking functions""" - - def __init__(self, tl_model, infer_type="default"): - self.tl_model = tl_model - self.infer_type = infer_type - - def predict(self, X): - import treelite - - if self.infer_type == "per_tree": - return treelite.gtil.predict_per_tree(self.tl_model, X) - return treelite.gtil.predict(self.tl_model, X) - - -def _build_gtil_classifier(m, data, args, tmpdir): - """Setup function for treelite classification benchmarking""" - import treelite - import xgboost as xgb - - max_depth = args["max_depth"] - num_rounds = args["num_rounds"] - infer_type = args.get("infer_type", "default") - n_feature = data[0].shape[1] - train_size = data[0].shape[0] - model_name = f"xgb_{max_depth}_{num_rounds}_{n_feature}_{train_size}.ubj" - model_path = os.path.join(tmpdir, model_name) - - bst = xgb.Booster() - bst.load_model(model_path) - tl_model = treelite.Model.from_xgboost(bst) - return GtilWrapper(tl_model, infer_type=infer_type) - - -def _treelite_fil_accuracy_score(y_true, y_pred): - """Function to get correct accuracy for FIL (returns class index)""" - # convert the input if necessary - if is_cuml_available(): - y_pred1 = ( - y_pred.copy_to_host() - if cuda.devicearray.is_cuda_ndarray(y_pred) - else y_pred - ) - y_true1 = ( - y_true.copy_to_host() - if cuda.devicearray.is_cuda_ndarray(y_true) - else y_true - ) - y_pred_binary = input_utils.convert_dtype(y_pred1 > 0.5, np.int32) - return cuml.metrics.accuracy_score(y_true1, y_pred_binary) - else: - # CPU-only fallback using sklearn - y_pred_binary = (np.asarray(y_pred) > 0.5).astype(np.int32) - return sklearn_metrics.accuracy_score( - np.asarray(y_true), y_pred_binary - ) - - def _build_mnmg_umap(m, data, args, tmpdir): """Build multi-node multi-GPU UMAP model. diff --git a/python/cuml/cuml/ensemble/randomforest_common.pyx b/python/cuml/cuml/ensemble/randomforest_common.pyx index a886aeb4b7..2bd13108a9 100644 --- a/python/cuml/cuml/ensemble/randomforest_common.pyx +++ b/python/cuml/cuml/ensemble/randomforest_common.pyx @@ -12,7 +12,7 @@ import cupy as cp import numpy as np import treelite.sklearn -from cuml.fil.fil import ForestInference +from cuml.fil.compat import ForestInference from cuml.internals.base import Base, get_handle from cuml.internals.interop import ( InteropMixin, @@ -26,6 +26,7 @@ from cuml.metrics import accuracy_score, r2_score from libc.stdint cimport uint64_t, uintptr_t from libcpp cimport bool from pylibraft.common.handle cimport handle_t +import nvforest from cuml.internals.logger cimport level_enum from cuml.internals.treelite cimport ( @@ -347,8 +348,8 @@ class BaseRandomForestModel(Base, InteropMixin): def __getstate__(self): state = self.__dict__.copy() - # FIL model isn't currently pickleable - state.pop("_fil_model", None) + # nvForest model isn't currently pickleable + state.pop("_nvforest_model", None) return state def __setstate__(self, state): @@ -374,13 +375,17 @@ class BaseRandomForestModel(Base, InteropMixin): self, layout="depth_first", default_chunk_size=None, align_bytes=None, ): """ - Create a Forest Inference (FIL) model from the trained cuML - Random Forest model. + Create a Forest Inference (FIL) model from the cuML model. + + .. deprecated:: 26.06 + + The ``as_fil`` method is deprecated and will be removed in 26.10. + Please use ``as_nvforest`` instead. Parameters ---------- layout : string (default = 'depth_first') - Specifies the in-memory layout of nodes in FIL forests. Options: + Specifies the in-memory layout of nodes in forests. Options: 'depth_first', 'layered', 'breadth_first'. default_chunk_size : int, optional (default = None) Determines how batches are further subdivided for parallel processing. @@ -399,7 +404,11 @@ class BaseRandomForestModel(Base, InteropMixin): inferencing on the random forest model. """ check_is_fitted(self) - + warnings.warn( + "as_fil() method is deprecated and will be removed in 26.10. " + "Use the as_nvforest() method instead.", + FutureWarning, + ) return ForestInference( verbose=self.verbose, output_type=self.output_type, @@ -409,6 +418,45 @@ class BaseRandomForestModel(Base, InteropMixin): default_chunk_size=default_chunk_size, align_bytes=align_bytes, ensure_all_finite=True, + _suppress_deprecation_warning=True, + ) + + def as_nvforest( + self, layout="depth_first", default_chunk_size=None, align_bytes=None, + ): + """ + Create a nvForest model from the cuML model. + + Parameters + ---------- + layout : string (default = 'depth_first') + Specifies the in-memory layout of nodes in forests. Options: + 'depth_first', 'layered', 'breadth_first'. + default_chunk_size : int, optional (default = None) + Determines how batches are further subdivided for parallel processing. + The optimal value depends on hardware, model, and batch size. + If None, will be automatically determined. + align_bytes : int, optional (default = None) + If specified, trees will be padded such that their in-memory size is + a multiple of this value. This can improve performance by guaranteeing + that memory reads from trees begin on a cache line boundary. + Typical values are 0 or 128 on GPU and 0 or 64 on CPU. + + Returns + ------- + nvforest_model : nvforest.ForestInference + A forest inference model which can be used to perform + inferencing on the random forest model. + """ + check_is_fitted(self) + + return nvforest.load_from_treelite_model( + tl_model=treelite.Model.deserialize_bytes(self._treelite_model_bytes), + device="gpu", + layout=layout, + default_chunk_size=default_chunk_size, + align_bytes=align_bytes, + handle=get_handle(), ) def _fit_forest(self, X, y): @@ -595,8 +643,8 @@ class BaseRandomForestModel(Base, InteropMixin): ) self.n_outputs_ = 1 self._treelite_model_bytes = (tl_bytes[:tl_bytes_len]) - # Ensure cached fil model is reset - self._fil_model = None + # Ensure cached nvforest model is reset + self._nvforest_model = None # Compute OOB score if requested if self.oob_score: @@ -605,7 +653,7 @@ class BaseRandomForestModel(Base, InteropMixin): self.feature_importances_ = feature_importances return self - def _get_inference_fil_model( + def _get_inference_nvforest_model( self, layout="depth_first", default_chunk_size=None, @@ -614,26 +662,26 @@ class BaseRandomForestModel(Base, InteropMixin): if ( layout == "depth_first" and default_chunk_size is None and align_bytes is None ): - # default parameters, get (or create) the cached fil model - if (fil_model := getattr(self, "_fil_model", None)) is None: - fil_model = self._fil_model = self.as_fil() + # default parameters, get (or create) the cached nvforest model + if (nvforest_model := getattr(self, "_nvforest_model", None)) is None: + nvforest_model = self._nvforest_model = self.as_nvforest() else: - fil_model = self.as_fil( + nvforest_model = self.as_nvforest( layout=layout, default_chunk_size=default_chunk_size, align_bytes=align_bytes, ) - return fil_model + return nvforest_model def _compute_oob_score(self, X, y, bootstrap_masks_cp): """ Compute OOB score using per-tree predictions and bootstrap masks. """ - # Get per-tree predictions using FIL - fil_model = self.as_fil() + # Get per-tree predictions using nvForest + nvforest_model = self._get_inference_nvforest_model() # Per tree predictions shape: (n_samples, n_trees) for regression # or (n_samples, n_trees, n_classes) for classification - per_tree_preds = fil_model.predict_per_tree(X) + per_tree_preds = nvforest_model.predict_per_tree(X) n_samples = X.shape[0] diff --git a/python/cuml/cuml/ensemble/randomforestclassifier.py b/python/cuml/cuml/ensemble/randomforestclassifier.py index b2a8517cbf..cf688fd9f6 100644 --- a/python/cuml/cuml/ensemble/randomforestclassifier.py +++ b/python/cuml/cuml/ensemble/randomforestclassifier.py @@ -11,7 +11,7 @@ from cuml.internals.array import CumlArray from cuml.internals.interop import UnsupportedOnGPU from cuml.internals.mixins import ClassifierMixin -from cuml.internals.validation import check_features, check_inputs +from cuml.internals.validation import check_inputs from cuml.metrics import accuracy_score @@ -281,15 +281,21 @@ def predict( ------- y : {} """ - fil = self._get_inference_fil_model( + nvforest_model = self._get_inference_nvforest_model( layout=layout, default_chunk_size=default_chunk_size, align_bytes=align_bytes, ) - check_features(self, X) - inds = fil.predict(X, threshold=threshold) - index = inds.index - inds = inds.to_output("cupy") + X_converted, index = check_inputs( + self, + X, + dtype=nvforest_model.forest.get_dtype(), + convert_dtype=convert_dtype, + order="C", + mem_type="device", + return_index=True, + ) + inds = nvforest_model.predict(X_converted, threshold=threshold) with cuml.internals.exit_internal_context(): output_type = self._get_output_type(X) return decode_labels( @@ -336,13 +342,21 @@ def predict_proba( ------- y : {} """ - fil = self._get_inference_fil_model( + nvforest_model = self._get_inference_nvforest_model( layout=layout, default_chunk_size=default_chunk_size, align_bytes=align_bytes, ) - check_features(self, X) - return fil.predict_proba(X) + X, index = check_inputs( + self, + X, + dtype=nvforest_model.forest.get_dtype(), + convert_dtype=convert_dtype, + order="C", + mem_type="device", + return_index=True, + ) + return CumlArray(nvforest_model.predict_proba(X), index=index) @insert_into_docstring( parameters=[("dense", "(n_samples, n_features)")], diff --git a/python/cuml/cuml/ensemble/randomforestregressor.py b/python/cuml/cuml/ensemble/randomforestregressor.py index ef31f4c216..026e5c17b6 100644 --- a/python/cuml/cuml/ensemble/randomforestregressor.py +++ b/python/cuml/cuml/ensemble/randomforestregressor.py @@ -7,7 +7,7 @@ from cuml.internals.array import CumlArray from cuml.internals.mixins import RegressorMixin from cuml.internals.outputs import reflect, run_in_internal_context -from cuml.internals.validation import check_features, check_inputs +from cuml.internals.validation import check_inputs from cuml.metrics import r2_score @@ -240,19 +240,27 @@ def predict( ------- y : {} """ - fil = self._get_inference_fil_model( + nvforest_model = self._get_inference_nvforest_model( layout=layout, default_chunk_size=default_chunk_size, align_bytes=align_bytes, ) - check_features(self, X) - preds = fil.predict(X) + X, index = check_inputs( + self, + X, + dtype=nvforest_model.forest.get_dtype(), + convert_dtype=convert_dtype, + order="C", + mem_type="device", + return_index=True, + ) + preds = nvforest_model.predict(X) # Reshape to 1D array if the output would be (n, 1) to match # the output shape behavior of scikit-learn. if len(preds.shape) == 2 and preds.shape[1] == 1: - preds = CumlArray(preds.to_output("cupy").reshape(-1)) - return preds + preds = preds.reshape(-1) + return CumlArray(preds, index=index) @nvtx.annotate( message="score RF-Regressor @randomforestclassifier.pyx", diff --git a/python/cuml/cuml/fil/CMakeLists.txt b/python/cuml/cuml/fil/CMakeLists.txt deleted file mode 100644 index a206a6514a..0000000000 --- a/python/cuml/cuml/fil/CMakeLists.txt +++ /dev/null @@ -1,19 +0,0 @@ -# ============================================================================= -# cmake-format: off -# SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# cmake-format: on -# ============================================================================= - -set(cython_sources "") -add_module_gpu_default( - "fil.pyx" ${fil_algo} ${randomforestclassifier_algo} ${randomforestregressor_algo} -) - -set(linked_libraries "${cuml_sg_libraries}" "${CUML_PYTHON_TREELITE_TARGET}") - -rapids_cython_create_modules( - CXX - SOURCE_FILES "${cython_sources}" - LINKED_LIBRARIES "${linked_libraries}" MODULE_PREFIX fil_ -) diff --git a/python/cuml/cuml/fil/__init__.py b/python/cuml/cuml/fil/__init__.py index b86cd3a346..2914986a67 100644 --- a/python/cuml/cuml/fil/__init__.py +++ b/python/cuml/cuml/fil/__init__.py @@ -1,9 +1,9 @@ # -# SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # -from cuml.fil.fil import ( +from cuml.fil.compat import ( ForestInference, get_fil_device_type, set_fil_device_type, diff --git a/python/cuml/cuml/fil/compat.py b/python/cuml/cuml/fil/compat.py new file mode 100644 index 0000000000..16665695eb --- /dev/null +++ b/python/cuml/cuml/fil/compat.py @@ -0,0 +1,514 @@ +# +# SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. +# SPDX-License-Identifier: Apache-2.0 +# + +""" +Compatibility shim for cuml.fil -> nvForest + +TODO(26.10): This module will be removed in 26.10. +""" + +import warnings + +import cupy as cp +import nvforest +import treelite + +from cuml.internals.array import CumlArray +from cuml.internals.base import Base, get_handle +from cuml.internals.device_type import DeviceType +from cuml.internals.global_settings import GlobalSettings +from cuml.internals.mixins import CMajorInputTagMixin +from cuml.internals.outputs import reflect +from cuml.internals.validation import check_array + + +def _is_nvforest_model_on_device(nvforest_model) -> bool: + return isinstance( + nvforest_model, + ( + nvforest.GPUForestInferenceClassifier, + nvforest.GPUForestInferenceRegressor, + ), + ) + + +class set_fil_device_type: + """Set the device type used by FIL. + + May optionally be used as a context-manager to set the device type only + within a context. + + This is deprecated and will be removed in 26.10. + + Parameters + ---------- + device_type : {'cpu', 'gpu'} + The device type to use. + + Examples + -------- + >>> from cuml.fil import set_fil_device_type # doctest: +SKIP + + Set the device type globally to use CPU. + + >>> set_fil_device_type("cpu") # doctest: +SKIP + + Set the device type globally to use GPU. + + >>> set_fil_device_type("gpu") # doctest: +SKIP + + Set the device type to use CPU within a context. + + >>> with set_fil_device_type("cpu"): # doctest: +SKIP + ... ... + """ + + def __init__(self, device_type): + warnings.warn( + "cuml.fil.set_fil_device_type is deprecated and will be removed in 26.10. ", + FutureWarning, + ) + device_type = DeviceType.from_str(device_type) + self._previous = GlobalSettings().fil_device_type + GlobalSettings().fil_device_type = device_type + + def __enter__(self): + return self + + def __exit__(self, *_): + GlobalSettings().fil_device_type = self._previous + + +def get_fil_device_type() -> DeviceType: + """Get the device type used by FIL.""" + return GlobalSettings().fil_device_type + + +class ForestInference(Base, CMajorInputTagMixin): + def __init__( + self, + *, + treelite_model=None, + output_type=None, + verbose=False, + is_classifier=False, + layout="depth_first", + default_chunk_size=None, + align_bytes=None, + precision="single", + device_id=None, + ensure_all_finite=False, + _suppress_deprecation_warning=False, + ): + super().__init__(verbose=verbose, output_type=output_type) + if not _suppress_deprecation_warning: + warnings.warn( + "cuml.fil.ForestInference is deprecated and will be removed in 26.10. " + "Use nvforest.load_model() or nvforest.load_from_sklearn() instead.", + FutureWarning, + stacklevel=2, + ) + if treelite_model is None: + self.model = None + elif isinstance(treelite_model, (treelite.Model, bytes)): + if isinstance(treelite_model, bytes): + treelite_model = treelite.Model.deserialize_bytes( + treelite_model + ) + self.model = nvforest.load_from_treelite_model( + tl_model=treelite_model, + device="gpu" + if get_fil_device_type() == DeviceType.device + else "cpu", + layout=layout, + default_chunk_size=default_chunk_size, + align_bytes=align_bytes, + precision=precision, + device_id=device_id, + handle=get_handle(), + ) + else: + raise ValueError( + f"Unrecognized type for treelite_model: {type(treelite_model)}" + ) + self.ensure_all_finite = ensure_all_finite + self._suppress_deprecation_warning = _suppress_deprecation_warning + + @property + def align_bytes(self): + return self.model.align_bytes if self.model else None + + @align_bytes.setter + def align_bytes(self, value): + raise NotImplementedError( + "Setter for align_bytes is no longer supported" + ) + + @property + def precision(self): + return self.model.precision if self.model else None + + @precision.setter + def precision(self, value): + raise NotImplementedError( + "Setter for precision is no longer supported" + ) + + @property + def is_classifier(self): + return self.model.is_classifier if self.model else None + + @is_classifier.setter + def is_classifier(self, value): + raise NotImplementedError( + "Setter for is_classifier is no longer supported" + ) + + @property + def device_id(self): + return self.model.device_id if self.model else None + + @device_id.setter + def device_id(self, value): + raise NotImplementedError( + "Setter for device_id is no longer supported" + ) + + @property + def treelite_model(self): + warnings.warn( + "Attribute treelite_model is no longer supported", + FutureWarning, + stacklevel=2, + ) + return None + + @treelite_model.setter + def treelite_model(self, value): + raise NotImplementedError( + "Setter for treelite_model is no longer supported" + ) + + @property + def layout(self): + return self.model.layout if self.model else None + + @layout.setter + def layout(self, value): + raise NotImplementedError("Setter for layout is no longer supported") + + def num_outputs(self): + return self.model.num_outputs if self.model else None + + def num_trees(self): + return self.model.num_trees if self.model else None + + @property + def default_chunk_size(self): + return self.model.default_chunk_size if self.model else None + + @classmethod + def load( + cls, + path, + *, + is_classifier=False, + precision="single", + model_type=None, + output_type=None, + verbose=False, + default_chunk_size=None, + align_bytes=None, + layout="depth_first", + device_id=0, + ): + warnings.warn( + "cuml.fil.ForestInference.load() is deprecated and will be removed in 26.10. " + "Use nvforest.load_model() instead.", + FutureWarning, + stacklevel=2, + ) + obj = cls( + output_type=output_type, + verbose=verbose, + _suppress_deprecation_warning=True, + ) + obj.model = nvforest.load_model( + model_file=path, + model_type=model_type, + device="gpu" + if get_fil_device_type() == DeviceType.device + else "cpu", + layout=layout, + default_chunk_size=default_chunk_size, + align_bytes=align_bytes, + precision=precision, + device_id=device_id, + handle=get_handle(), + ) + return obj + + @classmethod + def load_from_sklearn( + cls, + skl_model, + *, + is_classifier=False, + precision="single", + model_type=None, + output_type=None, + verbose=False, + default_chunk_size=None, + align_bytes=None, + layout="depth_first", + device_id=0, + ): + warnings.warn( + "cuml.fil.ForestInference.load_from_sklearn() is deprecated " + "and will be removed in 26.10. " + "Use nvforest.load_from_sklearn() instead.", + FutureWarning, + stacklevel=2, + ) + obj = cls( + output_type=output_type, + verbose=verbose, + _suppress_deprecation_warning=True, + ) + obj.model = nvforest.load_from_sklearn( + skl_model=skl_model, + device="gpu" + if get_fil_device_type() == DeviceType.device + else "cpu", + layout=layout, + default_chunk_size=default_chunk_size, + align_bytes=align_bytes, + precision=precision, + device_id=device_id, + handle=get_handle(), + ) + return obj + + @classmethod + def load_from_treelite_model( + cls, + tl_model, + *, + is_classifier=False, + precision="single", + model_type=None, + output_type=None, + verbose=False, + default_chunk_size=None, + align_bytes=None, + layout="depth_first", + device_id=0, + ): + warnings.warn( + "cuml.fil.ForestInference.load_from_treelite_model() is deprecated " + "and will be removed in 26.10. " + "Use nvforest.load_from_treelite_model() instead.", + FutureWarning, + stacklevel=2, + ) + obj = cls( + output_type=output_type, + verbose=verbose, + _suppress_deprecation_warning=True, + ) + obj.model = nvforest.load_from_treelite_model( + tl_model=tl_model, + device="gpu" + if get_fil_device_type() == DeviceType.device + else "cpu", + layout=layout, + default_chunk_size=default_chunk_size, + align_bytes=align_bytes, + precision=precision, + device_id=device_id, + handle=get_handle(), + ) + return obj + + def get_dtype(self): + if self.model is None: + raise RuntimeError("ForestInference not yet loaded") + return self.model.forest.get_dtype() + + @reflect + def predict_proba( + self, + X, + *, + preds=None, + chunk_size=None, + ) -> CumlArray: + if self.model is None: + raise RuntimeError("ForestInference not yet loaded") + if preds is not None: + raise NotImplementedError( + "Setting preds argument is no longer supported" + ) + if isinstance( + self.model, + ( + nvforest.GPUForestInferenceClassifier, + nvforest.CPUForestInferenceClassifier, + ), + ): + X, index = check_array( + X, + dtype=self.get_dtype(), + order="C", + mem_type="device" + if _is_nvforest_model_on_device(self.model) + else "host", + return_index=True, + ensure_all_finite=self.ensure_all_finite, + input_name="X", + ) + out = self.model.predict_proba(X, chunk_size=chunk_size) + mem_type = GlobalSettings().fil_memory_type.name + out = cp.asarray(out) if mem_type == "device" else cp.asnumpy(out) + return CumlArray(out, index=index) + raise RuntimeError("Must be a classifier to run predict_proba()") + + @reflect + def predict( + self, + X, + *, + preds=None, + chunk_size=None, + threshold=None, + ) -> CumlArray: + if self.model is None: + raise RuntimeError("ForestInference not yet loaded") + if preds is not None: + raise NotImplementedError( + "Setting preds argument is no longer supported" + ) + X, index = check_array( + X, + dtype=self.get_dtype(), + order="C", + mem_type="device" + if _is_nvforest_model_on_device(self.model) + else "host", + return_index=True, + ensure_all_finite=self.ensure_all_finite, + input_name="X", + ) + if isinstance( + self.model, + ( + nvforest.GPUForestInferenceClassifier, + nvforest.CPUForestInferenceClassifier, + ), + ): + out = self.model.predict( + X, chunk_size=chunk_size, threshold=threshold + ) + elif isinstance( + self.model, + ( + nvforest.GPUForestInferenceRegressor, + nvforest.CPUForestInferenceRegressor, + ), + ): + out = self.model.predict(X, chunk_size=chunk_size) + else: + raise NotImplementedError( + f"Unrecognized type for self.model: {type(self.model)}" + ) + mem_type = GlobalSettings().fil_memory_type.name + out = cp.asarray(out) if mem_type == "device" else cp.asnumpy(out) + return CumlArray(out, index=index) + + @reflect + def predict_per_tree(self, X, *, preds=None, chunk_size=None): + if self.model is None: + raise RuntimeError("ForestInference not yet loaded") + if preds is not None: + raise NotImplementedError( + "Setting preds argument is no longer supported" + ) + X, index = check_array( + X, + dtype=self.get_dtype(), + order="C", + mem_type="device" + if _is_nvforest_model_on_device(self.model) + else "host", + return_index=True, + ensure_all_finite=self.ensure_all_finite, + input_name="X", + ) + out = self.model.predict_per_tree(X, chunk_size=chunk_size) + mem_type = GlobalSettings().fil_memory_type.name + out = cp.asarray(out) if mem_type == "device" else cp.asnumpy(out) + return CumlArray(out, index=index) + + @reflect + def apply(self, X, *, preds=None, chunk_size=None): + if self.model is None: + raise RuntimeError("ForestInference not yet loaded") + if preds is not None: + raise NotImplementedError( + "Setting preds argument is no longer supported" + ) + X, index = check_array( + X, + dtype=self.get_dtype(), + order="C", + mem_type="device" + if _is_nvforest_model_on_device(self.model) + else "host", + return_index=True, + ensure_all_finite=self.ensure_all_finite, + input_name="X", + ) + out = self.model.apply(X, chunk_size=chunk_size) + mem_type = GlobalSettings().fil_memory_type.name + out = cp.asarray(out) if mem_type == "device" else cp.asnumpy(out) + return CumlArray(out, index=index) + + def optimize( + self, + *, + data=None, + batch_size=1024, + unique_batches=10, + timeout=0.2, + predict_method="predict", + max_chunk_size=None, + seed=0, + ): + if self.model is None: + raise RuntimeError("ForestInference not yet loaded") + return self.model.optimize( + data=data, + batch_size=batch_size, + unique_batches=unique_batches, + timeout=timeout, + predict_method=predict_method, + max_chunk_size=max_chunk_size, + seed=seed, + ) + + @classmethod + def _get_param_names(cls): + return [ + *super()._get_param_names(), + "treelite_model", + "is_classifier", + "layout", + "default_chunk_size", + "align_bytes", + "precision", + "device_id", + "ensure_all_finite", + "_suppress_deprecation_warning", + ] diff --git a/python/cuml/cuml/fil/detail/__init__.py b/python/cuml/cuml/fil/detail/__init__.py deleted file mode 100644 index b25a9d5a21..0000000000 --- a/python/cuml/cuml/fil/detail/__init__.py +++ /dev/null @@ -1,4 +0,0 @@ -# -# SPDX-FileCopyrightText: Copyright (c) 2023, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# diff --git a/python/cuml/cuml/fil/detail/raft_proto/__init__.py b/python/cuml/cuml/fil/detail/raft_proto/__init__.py deleted file mode 100644 index b25a9d5a21..0000000000 --- a/python/cuml/cuml/fil/detail/raft_proto/__init__.py +++ /dev/null @@ -1,4 +0,0 @@ -# -# SPDX-FileCopyrightText: Copyright (c) 2023, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# diff --git a/python/cuml/cuml/fil/detail/raft_proto/cuda_stream.pxd b/python/cuml/cuml/fil/detail/raft_proto/cuda_stream.pxd deleted file mode 100644 index fcd52b70d4..0000000000 --- a/python/cuml/cuml/fil/detail/raft_proto/cuda_stream.pxd +++ /dev/null @@ -1,7 +0,0 @@ -# -# SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# -cdef extern from "cuml/fil/detail/raft_proto/cuda_stream.hpp" namespace "raft_proto" nogil: - cdef cppclass cuda_stream: - pass diff --git a/python/cuml/cuml/fil/detail/raft_proto/device_type.pxd b/python/cuml/cuml/fil/detail/raft_proto/device_type.pxd deleted file mode 100644 index 509eaff071..0000000000 --- a/python/cuml/cuml/fil/detail/raft_proto/device_type.pxd +++ /dev/null @@ -1,8 +0,0 @@ -# -# SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# -cdef extern from "cuml/fil/detail/raft_proto/device_type.hpp" namespace "raft_proto" nogil: - cdef enum device_type: - cpu "raft_proto::device_type::cpu", - gpu "raft_proto::device_type::gpu" diff --git a/python/cuml/cuml/fil/detail/raft_proto/handle.pxd b/python/cuml/cuml/fil/detail/raft_proto/handle.pxd deleted file mode 100644 index d26cafe969..0000000000 --- a/python/cuml/cuml/fil/detail/raft_proto/handle.pxd +++ /dev/null @@ -1,19 +0,0 @@ -# -# SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# - -from pylibraft.common.handle cimport handle_t as raft_handle_t - -from cuml.fil.detail.raft_proto.cuda_stream cimport ( - cuda_stream as raft_proto_stream_t, -) - - -cdef extern from "cuml/fil/detail/raft_proto/handle.hpp" namespace "raft_proto" nogil: - cdef cppclass handle_t: - handle_t() except + - handle_t(const raft_handle_t* handle_ptr) except + - handle_t(const raft_handle_t& handle) except + - raft_proto_stream_t get_next_usable_stream() except + - void synchronize() except+ diff --git a/python/cuml/cuml/fil/detail/raft_proto/optional.pxd b/python/cuml/cuml/fil/detail/raft_proto/optional.pxd deleted file mode 100644 index 54dcac4dc2..0000000000 --- a/python/cuml/cuml/fil/detail/raft_proto/optional.pxd +++ /dev/null @@ -1,43 +0,0 @@ -# -# SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# -# The following is taken from -# https://github.com/cython/cython/blob/master/Cython/Includes/libcpp/optional.pxd, -# which provides a binding for std::optional in Cython 3.0 - -from libcpp cimport bool - - -cdef extern from "" namespace "std" nogil: - cdef cppclass nullopt_t: - nullopt_t() - - cdef nullopt_t nullopt - - cdef cppclass optional[T]: - ctypedef T value_type - optional() - optional(nullopt_t) - optional(optional&) except + - optional(T&) except + - bool has_value() - T& value() - T& value_or[U](U& default_value) - void swap(optional&) - void reset() - T& emplace(...) - T& operator*() - # T* operator->() # Not Supported - optional& operator=(optional&) - optional& operator=[U](U&) - bool operator bool() - bool operator!() - bool operator==[U](optional&, U&) - bool operator!=[U](optional&, U&) - bool operator<[U](optional&, U&) - bool operator>[U](optional&, U&) - bool operator<=[U](optional&, U&) - bool operator>=[U](optional&, U&) - - optional[T] make_optional[T](...) except + diff --git a/python/cuml/cuml/fil/fil.pyx b/python/cuml/cuml/fil/fil.pyx deleted file mode 100644 index af12a3736f..0000000000 --- a/python/cuml/cuml/fil/fil.pyx +++ /dev/null @@ -1,1371 +0,0 @@ -# -# SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# -import itertools -import pathlib -from time import perf_counter - -import cupy as cp -import numpy as np -import treelite.sklearn -from cuda.bindings import runtime - -import cuml.internals.nvtx as nvtx -from cuml.internals.array import CumlArray -from cuml.internals.base import Base, get_handle -from cuml.internals.device_type import DeviceType, DeviceTypeError -from cuml.internals.global_settings import GlobalSettings -from cuml.internals.mem_type import MemoryType -from cuml.internals.mixins import CMajorInputTagMixin -from cuml.internals.outputs import reflect -from cuml.internals.treelite import safe_treelite_call -from cuml.internals.validation import check_array - -from libc.stdint cimport uint32_t, uintptr_t -from libcpp cimport bool -from pylibraft.common.handle cimport handle_t as raft_handle_t - -from cuml.fil.detail.raft_proto.cuda_stream cimport ( - cuda_stream as raft_proto_stream_t, -) -from cuml.fil.detail.raft_proto.device_type cimport ( - device_type as raft_proto_device_t, -) -from cuml.fil.detail.raft_proto.handle cimport handle_t as raft_proto_handle_t -from cuml.fil.detail.raft_proto.optional cimport nullopt, optional -from cuml.fil.infer_kind cimport infer_kind -from cuml.fil.postprocessing cimport element_op, row_op -from cuml.fil.tree_layout cimport tree_layout as fil_tree_layout -from cuml.internals.treelite cimport ( - TreeliteDeserializeModelFromBytes, - TreeliteFreeModel, - TreeliteModelHandle, -) - - -cdef extern from "cuml/fil/forest_model.hpp" namespace "ML::fil" nogil: - cdef cppclass forest_model: - void predict[io_t]( - const raft_proto_handle_t&, - io_t*, - io_t*, - size_t, - raft_proto_device_t, - raft_proto_device_t, - infer_kind, - optional[uint32_t] - ) except + - - bool is_double_precision() except + - size_t num_features() except + - size_t num_outputs() except + - size_t num_trees() except + - bool has_vector_leaves() except + - row_op row_postprocessing() except + - element_op elem_postprocessing() except + - -cdef extern from "cuml/fil/treelite_importer.hpp" namespace "ML::fil" nogil: - forest_model import_from_treelite_handle( - TreeliteModelHandle, - fil_tree_layout, - uint32_t, - optional[bool], - raft_proto_device_t, - int, - raft_proto_stream_t - ) except + - - -class set_fil_device_type: - """Set the device type used by FIL. - - May optionally be used as a context-manager to set the device type only - within a context. - - Parameters - ---------- - device_type : {'cpu', 'gpu'} - The device type to use. - - Examples - -------- - >>> from cuml.fil import set_fil_device_type # doctest: +SKIP - - Set the device type globally to use CPU. - - >>> set_fil_device_type("cpu") # doctest: +SKIP - - Set the device type globally to use GPU. - - >>> set_fil_device_type("gpu") # doctest: +SKIP - - Set the device type to use CPU within a context. - - >>> with set_fil_device_type("cpu"): # doctest: +SKIP - ... ... - """ - def __init__(self, device_type): - device_type = DeviceType.from_str(device_type) - self._previous = GlobalSettings().fil_device_type - GlobalSettings().fil_device_type = device_type - - def __enter__(self): - return self - - def __exit__(self, *_): - GlobalSettings().fil_device_type = self._previous - - -def get_fil_device_type() -> DeviceType: - """Get the device type used by FIL.""" - return GlobalSettings().fil_device_type - - -cdef raft_proto_device_t get_fil_raft_proto_device_type(arr): - """Get the current FIL device type as a raft_proto_device_t""" - if isinstance(arr, cp.ndarray): - return raft_proto_device_t.gpu - elif isinstance(arr, np.ndarray): - return raft_proto_device_t.cpu - else: - if arr.mem_type is MemoryType.device: - return raft_proto_device_t.gpu - else: - return raft_proto_device_t.cpu - - -cdef class ForestInference_impl(): - cdef forest_model model - cdef raft_proto_handle_t raft_proto_handle - cdef object raft_handle - cdef object ensure_all_finite - - def __cinit__( - self, - tl_model_bytes, - *, - layout='depth_first', - align_bytes=0, - use_double_precision=None, - mem_type=None, - device_id=None, - ensure_all_finite=False, - ): - # Store reference to RAFT handle to control lifetime, since raft_proto - # handle keeps a pointer to it - self.raft_handle = get_handle() - self.raft_proto_handle = raft_proto_handle_t( - self.raft_handle.getHandle() - ) - self.ensure_all_finite = ensure_all_finite - if mem_type is None: - mem_type = GlobalSettings().fil_memory_type - else: - mem_type = MemoryType.from_str(mem_type) - - cdef optional[bool] use_double_precision_c - cdef bool use_double_precision_bool - if use_double_precision is None: - use_double_precision_c = nullopt - else: - use_double_precision_bool = use_double_precision - use_double_precision_c = use_double_precision_bool - - cdef TreeliteModelHandle tl_handle = NULL - safe_treelite_call( - TreeliteDeserializeModelFromBytes( - tl_model_bytes, len(tl_model_bytes), &tl_handle), - "Failed to load Treelite model from bytes:" - ) - - cdef raft_proto_device_t dev_type - if mem_type is MemoryType.device: - dev_type = raft_proto_device_t.gpu - else: - dev_type = raft_proto_device_t.cpu - cdef fil_tree_layout tree_layout - if layout.lower() == "depth_first": - tree_layout = fil_tree_layout.depth_first - elif layout.lower() == "breadth_first": - tree_layout = fil_tree_layout.breadth_first - elif layout.lower() == "layered": - tree_layout = fil_tree_layout.layered_children_together - else: - raise RuntimeError(f"Unrecognized tree layout {layout}") - - # Use assertion here, since device_id being None would indicate - # a bug, not a user error. The outer ForestInference object - # should set an integer device_id before passing it to - # ForestInference_impl. - assert device_id is not None, ( - "device_id should be set before building ForestInference_impl" - ) - - self.model = import_from_treelite_handle( - tl_handle, - tree_layout, - align_bytes, - use_double_precision_c, - dev_type, - device_id, - self.raft_proto_handle.get_next_usable_stream() - ) - - safe_treelite_call( - TreeliteFreeModel(tl_handle), - "Failed to free Treelite model:" - ) - - def get_dtype(self): - return [np.float32, np.float64][self.model.is_double_precision()] - - def num_features(self): - return self.model.num_features() - - def num_outputs(self): - return self.model.num_outputs() - - def num_trees(self): - return self.model.num_trees() - - def row_postprocessing(self): - enum_val = self.model.row_postprocessing() - if enum_val == row_op.row_disable: - return "disable" - elif enum_val == row_op.softmax: - return "softmax" - elif enum_val == row_op.max_index: - return "max_index" - - def elem_postprocessing(self): - enum_val = self.model.elem_postprocessing() - if enum_val == element_op.elem_disable: - return "disable" - elif enum_val == element_op.signed_square: - return "signed_square" - elif enum_val == element_op.hinge: - return "hinge" - elif enum_val == element_op.sigmoid: - return "sigmoid" - elif enum_val == element_op.exponential: - return "exponential" - elif enum_val == element_op.logarithm_one_plus_exp: - return "logarithm_one_plus_exp" - - def _predict(self, X, *, predict_type="default", preds=None, chunk_size=None): - model_dtype = self.get_dtype() - mem_type = GlobalSettings().fil_memory_type - - X, index = check_array( - X, - dtype=model_dtype, - order="C", - mem_type=mem_type.name, - return_index=True, - ensure_all_finite=self.ensure_all_finite, - input_name="X", - ) - n_rows = X.shape[0] - - cdef raft_proto_device_t in_dev = get_fil_raft_proto_device_type(X) - cdef uintptr_t in_ptr = ( - X.data.ptr if isinstance(X, cp.ndarray) else X.ctypes.data - ) - - cdef uintptr_t out_ptr - cdef infer_kind infer_type_enum - if predict_type == "default": - infer_type_enum = infer_kind.default_kind - output_shape = (n_rows, self.model.num_outputs()) - elif predict_type == "per_tree": - infer_type_enum = infer_kind.per_tree - if self.model.has_vector_leaves(): - output_shape = (n_rows, self.model.num_trees(), self.model.num_outputs()) - else: - output_shape = (n_rows, self.model.num_trees()) - elif predict_type == "leaf_id": - infer_type_enum = infer_kind.leaf_id - output_shape = (n_rows, self.model.num_trees()) - else: - raise ValueError(f"Unrecognized predict_type: {predict_type}") - if preds is None: - preds = CumlArray.empty( - output_shape, - model_dtype, - order='C', - index=index, - mem_type=mem_type, - ) - else: - # TODO(wphicks): Handle incorrect dtype/device/layout in C++ - if preds.shape != output_shape: - raise ValueError(f"If supplied, preds argument must have shape {output_shape}") - preds.index = index - cdef raft_proto_device_t out_dev - out_dev = get_fil_raft_proto_device_type(preds) - out_ptr = preds.ptr - - cdef optional[uint32_t] chunk_specification - if chunk_size is None: - chunk_specification = nullopt - else: - chunk_specification = chunk_size - - if model_dtype == np.float32: - self.model.predict[float]( - self.raft_proto_handle, - out_ptr, - in_ptr, - n_rows, - out_dev, - in_dev, - infer_type_enum, - chunk_specification - ) - else: - self.model.predict[double]( - self.raft_proto_handle, - out_ptr, - in_ptr, - n_rows, - out_dev, - in_dev, - infer_type_enum, - chunk_specification - ) - - if get_fil_device_type() is DeviceType.device: - self.raft_proto_handle.synchronize() - return preds - - def predict( - self, - X, - *, - predict_type="default", - preds=None, - chunk_size=None, - ): - return self._predict( - X, - predict_type=predict_type, - preds=preds, - chunk_size=chunk_size, - ) - - -class _AutoIterations: - """Used to generate sequence of iterations (1, 2, 5, 10, 20, 50...) during - FIL optimization""" - - def __init__(self): - self.invocations = 0 - self.sequence = (1, 2, 5) - - def next(self): - result = ( - (10 ** ( - self.invocations // len(self.sequence) - )) * self.sequence[self.invocations % len(self.sequence)] - ) - self.invocations += 1 - return result - - -class ForestInference(Base, CMajorInputTagMixin): - """ - ForestInference provides accelerated inference for forest models on both - CPU and GPU. - - **Performance Tuning** - FIL offers a number of hyperparameters that can be tuned to obtain optimal - performance for a given model, hardware, and batch size. The easiest way to - optimize these parameters is using the automated `.optimize` method, which - will find the optimum for an indicated batch size. For some use cases, - manual adjustment of these parameters is preferred, so available - performance hyperparameters are described in detail below. - - To obtain optimal performance with this implementation of FIL, the single - most important value is the `chunk_size` parameter passed to the predict - method. Essentially, `chunk_size` determines how many rows to evaluate - together at once from a single batch. Larger values reduce global memory - accesses on GPU and cache misses on CPU, but smaller values allow for - finer-grained parallelism, improving usage of available processing power. - The optimal value for this parameter is hard to predict a priori, but in - general larger batch sizes benefit from larger chunk sizes and smaller - batch sizes benefit from smaller chunk sizes. Having a chunk size larger - than the batch size is never optimal. - - To determine the optimal chunk size on GPU, test powers of 2 from 1 to - 32. Values above 32 and values which are not powers of 2 are not supported. - - To determine the optimal chunk size on CPU, test powers of 2 from 1 to - 512. Values above 512 are supported, but RAPIDS developers have not yet - seen a case where they yield improved performance. - - After chunk size, the most important performance parameter is `layout`, - also described below. Testing available layouts is recommended to optimize - performance, but the impact is likely to be substantially less than - optimizing `chunk_size`. There is no universal rule for predicting which - layout will produce the best performance. On both GPU and CPU, the - `depth_first` layout can improve performance by increasing cache hits - during tree traversal. This tends to be the strongest effect for most use - cases, so `depth_first` is used as the default value. - - `align_bytes` is the final performance parameter. This parameter allows - trees to be padded with empty nodes until their total in-memory size is a - multiple of the given value. In general, if a non-default value is used, it - should either be 0 or the cache line byte size for the device being used - for execution (64 for CPU or 128 for GPU). If left unpadded, forest data - remains more compact in memory, which can improve the frequency of cache - hits. On the other hand, padding to the size of the cache line ensures that - trees begin on cache line boundaries. It is difficult to predict for any - given model which effect will be the greater determinant of performance. If - left at the default value of `None`, trees will be unpadded for GPU - execution and padded to 64 bytes for CPU execution. This value has no - effect for the `layered` layout, since trees in this layout overlap in - memory. - - Parameters - ---------- - treelite_model : treelite.Model - The model to be used for inference. This can be trained with XGBoost, - LightGBM, cuML, Scikit-Learn, or any other forest model framework - so long as it can be loaded into a treelite.Model object (See - https://treelite.readthedocs.io/en/latest/treelite-api.html). - output_type : {'input', 'array', 'dataframe', 'series', 'df_obj', \ - 'numba', 'cupy', 'numpy', 'cudf', 'pandas'}, default=None - Return results and set estimator attributes to the indicated output - type. If None, the output type set at the module level - (`cuml.global_settings.output_type`) will be used. See - :ref:`output-data-type-configuration` for more info. - verbose : int or boolean, default=False - Sets logging level. It must be one of `cuml.common.logger.level_*`. - See :ref:`verbosity-levels` for more info. - layout : {'breadth_first', 'depth_first', 'layered'}, default='depth_first' - The in-memory layout to be used during inference for nodes of the - forest model. This parameter is available purely for runtime - optimization. For performance-critical applications, it is - recommended that each layout be tested with realistic batch sizes to - determine the optimal value. - align_bytes : int or None, default=None - Pad each tree with empty nodes until its in-memory size is a multiple - of the given value. If None, use 0 for GPU and 64 for CPU. - precision : {'single', 'double', None}, default='single' - Use the given floating point precision for evaluating the model. If - None, use the native precision of the model. Note that - single-precision execution is substantially faster than - double-precision execution, so double-precision is recommended - only for models trained and double precision and when exact - conformance between results from FIL and the original training - framework is of paramount importance. - device_id : int or None, default=None - For GPU execution, the device on which to load and execute this - model. If set to None, use the currently active device. - For CPU execution, this value is currently ignored. - ensure_all_finite : bool or 'allow-nan', default=False - If True, an error will be raised if non-finite values are found in the - input. If 'allow-nan', an error will be raised if infinite values are - found (but not for NaN). If False then ``check_all_finite`` is skipped. - """ - - def _reload_model(self): - """Reload model on any device (CPU/GPU) where model has already been - loaded""" - if hasattr(self, '_gpu_forest'): - with set_fil_device_type('gpu'): - self._load_to_fil(device_id=self.device_id) - if hasattr(self, '_cpu_forest'): - with set_fil_device_type('cpu'): - self._load_to_fil(device_id=self.device_id) - - @staticmethod - def _get_default_align_bytes(): - if get_fil_device_type() is DeviceType.host: - return 64 - else: - return 0 - - @property - def align_bytes(self): - try: - return self._align_bytes_ - except AttributeError: - return self._get_default_align_bytes() - - @align_bytes.setter - def align_bytes(self, value): - try: - old_value = self._align_bytes_ - except AttributeError: - old_value = None - if value is None: - if old_value is not None: - del self._align_bytes_ - self._reload_model() - else: - self._align_bytes_ = value - if old_value is None or value != old_value: - self._reload_model() - - @property - def precision(self): - try: - use_double_precision = \ - self._use_double_precision_ - except AttributeError: - self._use_double_precision_ = False - use_double_precision = \ - self._use_double_precision_ - if use_double_precision is None: - return 'native' - elif use_double_precision: - return 'double' - else: - return 'single' - - @precision.setter - def precision(self, value): - try: - old_value = self._use_double_precision_ - except AttributeError: - self._use_double_precision_ = False - old_value = self._use_double_precision_ - if value in ('native', None): - self._use_double_precision_ = None - elif value in ('double', 'float64'): - self._use_double_precision_ = True - else: - self._use_double_precision_ = False - if old_value != self._use_double_precision_: - self._reload_model() - - @property - def is_classifier(self): - try: - return self._is_classifier_ - except AttributeError: - self._is_classifier_ = False - return self._is_classifier_ - - @is_classifier.setter - def is_classifier(self, value): - if not hasattr(self, '_is_classifier_'): - self._is_classifier_ = value - elif value is not None: - self._is_classifier_ = value - - @property - def device_id(self): - try: - return self._device_id_ - except AttributeError: - self._device_id_ = None - return self._device_id_ - - @device_id.setter - def device_id(self, value): - try: - old_value = self.device_id - except AttributeError: - old_value = None - self._device_id_ = value - if ( - self.treelite_model is not None - and self.device_id != old_value - and hasattr(self, '_gpu_forest') - ): - self._load_to_fil(device_id=self.device_id) - - @property - def treelite_model(self): - try: - return self._treelite_model_ - except AttributeError: - return None - - @treelite_model.setter - def treelite_model(self, value): - if value is not None: - self._treelite_model_ = value - self._reload_model() - - @property - def layout(self): - try: - return self._layout_ - except AttributeError: - self._layout_ = 'depth_first' - return self._layout_ - - @layout.setter - def layout(self, value): - try: - old_value = self._layout_ - except AttributeError: - old_value = None - if value is not None: - self._layout_ = value - if old_value != value: - self._reload_model() - - def __init__( - self, - *, - treelite_model=None, - output_type=None, - verbose=False, - is_classifier=False, - layout='depth_first', - default_chunk_size=None, - align_bytes=None, - precision='single', - device_id=None, - ensure_all_finite=False, - ): - super().__init__(verbose=verbose, output_type=output_type) - self.is_classifier = is_classifier - self.default_chunk_size = default_chunk_size - self.align_bytes = align_bytes - self.layout = layout - self.precision = precision - self.device_id = device_id - self.ensure_all_finite = ensure_all_finite - self.treelite_model = treelite_model - self._load_to_fil(device_id=self.device_id) - - def _load_to_fil(self, mem_type=None, device_id=None): - if mem_type is None: - mem_type = GlobalSettings().fil_memory_type - else: - mem_type = MemoryType.from_str(mem_type) - - if device_id is None: - # If no device ID is explicitly given, use the currently - # active device - status, current_device_id = runtime.cudaGetDevice() - if status != runtime.cudaError_t.cudaSuccess: - _, name = runtime.cudaGetErrorName(status) - _, msg = runtime.cudaGetErrorString(status) - raise RuntimeError(f"Failed to run cudaGetDevice(). {name}: {msg}") - device_id = current_device_id - - if mem_type is MemoryType.device: - self.device_id = device_id - - if self.treelite_model is not None: - if isinstance(self.treelite_model, treelite.Model): - treelite_model_bytes = self.treelite_model.serialize_bytes() - elif isinstance(self.treelite_model, bytes): - treelite_model_bytes = self.treelite_model - else: - raise ValueError("treelite_model should be either treelite.Model or bytes") - impl = ForestInference_impl( - treelite_model_bytes, - layout=self.layout, - align_bytes=self.align_bytes, - use_double_precision=self._use_double_precision_, - mem_type=mem_type, - device_id=self.device_id, - ensure_all_finite=self.ensure_all_finite - ) - - if mem_type is MemoryType.device: - self._gpu_forest = impl - else: - self._cpu_forest = impl - - @property - def gpu_forest(self): - """The underlying FIL forest model loaded in GPU-accessible memory""" - try: - return self._gpu_forest - except AttributeError: - self._load_to_fil(mem_type=MemoryType.device) - return self._gpu_forest - - @property - def cpu_forest(self): - """The underlying FIL forest model loaded in CPU-accessible memory""" - try: - return self._cpu_forest - except AttributeError: - self._load_to_fil(mem_type=MemoryType.host) - return self._cpu_forest - - @property - def forest(self): - """The underlying FIL forest model loaded in memory compatible with the - current global device_type setting""" - device_type = get_fil_device_type() - if device_type is DeviceType.device: - return self.gpu_forest - elif device_type is DeviceType.host: - return self.cpu_forest - else: - raise DeviceTypeError("Unsupported device type for FIL") - - def num_outputs(self): - return self.forest.num_outputs() - - def num_trees(self): - return self.forest.num_trees() - - @classmethod - def load( - cls, - path, - *, - is_classifier=False, - precision='single', - model_type=None, - output_type=None, - verbose=False, - default_chunk_size=None, - align_bytes=None, - layout='depth_first', - device_id=0, - ): - """Load a model into FIL from a serialized model file. - - Parameters - ---------- - path : str - The path to the serialized model file. This can be an XGBoost - binary or JSON file, a LightGBM text file, or a Treelite checkpoint - file. If the model_type parameter is not passed, an attempt will be - made to load the file based on its extension. - is_classifier : boolean, default=False - True for classification models, False for regressors - precision : {'single', 'double', None}, default='single' - Use the given floating point precision for evaluating the model. If - None, use the native precision of the model. Note that - single-precision execution is substantially faster than - double-precision execution, so double-precision is recommended - only for models trained and double precision and when exact - conformance between results from FIL and the original training - framework is of paramount importance. - model_type : {'xgboost_ubj', 'xgboost_json', 'xgboost', 'lightgbm', - 'treelite_checkpoint', None }, default=None - The serialization format for the model file. If None, a best-effort - guess will be made based on the file extension. - output_type : {'input', 'array', 'dataframe', 'series', 'df_obj', \ - 'numba', 'cupy', 'numpy', 'cudf', 'pandas'}, default=None - Return results and set estimator attributes to the indicated output - type. If None, the output type set at the module level - (`cuml.global_settings.output_type`) will be used. See - :ref:`output-data-type-configuration` for more info. - verbose : int or boolean, default=False - Sets logging level. It must be one of `cuml.common.logger.level_*`. - See :ref:`verbosity-levels` for more info. - default_chunk_size : int or None, default=None - If set, predict calls without a specified chunk size will use - this default value. - align_bytes : int or None, default=None - Pad each tree with empty nodes until its in-memory size is a multiple - of the given value. If None, use 0 for GPU and 64 for CPU. - layout : {'breadth_first', 'depth_first', 'layered'}, default='depth_first' - The in-memory layout to be used during inference for nodes of the - forest model. This parameter is available purely for runtime - optimization. For performance-critical applications, it is - recommended that available layouts be tested with realistic batch - sizes to determine the optimal value. - device_id : int, default=0 - For GPU execution, the device on which to load and execute this - model. For CPU execution, this value is currently ignored. - """ - if model_type is None: - extension = pathlib.Path(path).suffix - if extension == '.json': - model_type = 'xgboost_json' - elif extension == '.ubj': - model_type = 'xgboost_ubj' - elif extension == '.model': - model_type = 'xgboost' - elif extension == '.txt': - model_type = 'lightgbm' - else: - model_type = 'treelite_checkpoint' - if model_type == "treelite_checkpoint": - tl_model = treelite.frontend.Model.deserialize(path) - elif model_type == "xgboost_ubj": - tl_model = treelite.frontend.load_xgboost_model(path, format_choice="ubjson") - elif model_type == "xgboost_json": - tl_model = treelite.frontend.load_xgboost_model(path, format_choice="json") - elif model_type == "xgboost": - tl_model = treelite.frontend.load_xgboost_model_legacy_binary(path) - elif model_type == "lightgbm": - tl_model = treelite.frontend.load_lightgbm_model(path) - else: - raise ValueError(f"Unknown model type: {model_type}") - return cls( - treelite_model=tl_model, - output_type=output_type, - verbose=verbose, - is_classifier=is_classifier, - default_chunk_size=default_chunk_size, - align_bytes=align_bytes, - layout=layout, - precision=precision, - device_id=device_id - ) - - @classmethod - def load_from_sklearn( - cls, - skl_model, - *, - is_classifier=False, - precision='single', - model_type=None, - output_type=None, - verbose=False, - default_chunk_size=None, - align_bytes=None, - layout='depth_first', - device_id=0, - ): - """Load a Scikit-Learn forest model to FIL - - Parameters - ---------- - skl_model - The Scikit-Learn forest model to load. - is_classifier : boolean, default=False - True for classification models, False for regressors - precision : {'single', 'double', None}, default='single' - Use the given floating point precision for evaluating the model. If - None, use the native precision of the model. Note that - single-precision execution is substantially faster than - double-precision execution, so double-precision is recommended - only for models trained and double precision and when exact - conformance between results from FIL and the original training - framework is of paramount importance. - model_type : {'xgboost', 'xgboost_json', 'lightgbm', - 'treelite_checkpoint', None }, default=None - The serialization format for the model file. If None, a best-effort - guess will be made based on the file extension. - output_type : {'input', 'array', 'dataframe', 'series', 'df_obj', \ - 'numba', 'cupy', 'numpy', 'cudf', 'pandas'}, default=None - Return results and set estimator attributes to the indicated output - type. If None, the output type set at the module level - (`cuml.global_settings.output_type`) will be used. See - :ref:`output-data-type-configuration` for more info. - verbose : int or boolean, default=False - Sets logging level. It must be one of `cuml.common.logger.level_*`. - See :ref:`verbosity-levels` for more info. - default_chunk_size : int or None, default=None - If set, predict calls without a specified chunk size will use - this default value. - align_bytes : int or None, default=None - Pad each tree with empty nodes until its in-memory size is a multiple - of the given value. If None, use 0 for GPU and 64 for CPU. - layout : {'breadth_first', 'depth_first', 'layered'}, default='depth_first' - The in-memory layout to be used during inference for nodes of the - forest model. This parameter is available purely for runtime - optimization. For performance-critical applications, it is - recommended that available layouts be tested with realistic batch - sizes to determine the optimal value. - mem_type : {'device', 'host', None}, default='single' - The memory type to use for initially loading the model. If None, - the current global memory type setting will be used. If the model - is loaded with one memory type and inference is later requested - with an incompatible device (e.g. device memory and CPU execution), - the model will be lazily loaded to the correct location at that - time. In general, it should not be necessary to set this parameter - directly (rely instead on the `set_fil_device_type` context manager), - but it can be a useful convenience for some hyperoptimization - pipelines. - device_id : int, default=0 - For GPU execution, the device on which to load and execute this - model. For CPU execution, this value is currently ignored. - """ - tl_model = treelite.sklearn.import_model(skl_model) - result = cls( - treelite_model=tl_model, - output_type=output_type, - verbose=verbose, - is_classifier=is_classifier, - default_chunk_size=default_chunk_size, - align_bytes=align_bytes, - layout=layout, - precision=precision, - device_id=device_id - ) - return result - - @classmethod - def load_from_treelite_model( - cls, - tl_model, - *, - is_classifier=False, - precision='single', - model_type=None, - output_type=None, - verbose=False, - default_chunk_size=None, - align_bytes=None, - layout='depth_first', - device_id=0, - ): - """Load a Treelite model to FIL - - Parameters - ---------- - tl_model : treelite.Model - The Treelite model to load. - is_classifier : boolean, default=False - True for classification models, False for regressors - precision : {'single', 'double', None}, default='single' - Use the given floating point precision for evaluating the model. If - None, use the native precision of the model. Note that - single-precision execution is substantially faster than - double-precision execution, so double-precision is recommended - only for models trained and double precision and when exact - conformance between results from FIL and the original training - framework is of paramount importance. - model_type : {'xgboost', 'xgboost_json', 'lightgbm', - 'treelite_checkpoint', None }, default=None - The serialization format for the model file. If None, a best-effort - guess will be made based on the file extension. - output_type : {'input', 'array', 'dataframe', 'series', 'df_obj', \ - 'numba', 'cupy', 'numpy', 'cudf', 'pandas'}, default=None - Return results and set estimator attributes to the indicated output - type. If None, the output type set at the module level - (`cuml.global_settings.output_type`) will be used. See - :ref:`output-data-type-configuration` for more info. - verbose : int or boolean, default=False - Sets logging level. It must be one of `cuml.common.logger.level_*`. - See :ref:`verbosity-levels` for more info. - default_chunk_size : int or None, default=None - If set, predict calls without a specified chunk size will use - this default value. - align_bytes : int or None, default=None - Pad each tree with empty nodes until its in-memory size is a multiple - of the given value. If None, use 0 for GPU and 64 for CPU. - layout : {'breadth_first', 'depth_first', 'layered'}, default='depth_first' - The in-memory layout to be used during inference for nodes of the - forest model. This parameter is available purely for runtime - optimization. For performance-critical applications, it is - recommended that available layouts be tested with realistic batch - sizes to determine the optimal value. - mem_type : {'device', 'host', None}, default='single' - The memory type to use for initially loading the model. If None, - the current global memory type setting will be used. If the model - is loaded with one memory type and inference is later requested - with an incompatible device (e.g. device memory and CPU execution), - the model will be lazily loaded to the correct location at that - time. In general, it should not be necessary to set this parameter - directly (rely instead on the `set_fil_device_type` context - manager), but it can be a useful convenience for some - hyperoptimization pipelines. - device_id : int, default=0 - For GPU execution, the device on which to load and execute this - model. For CPU execution, this value is currently ignored. - """ - return cls( - treelite_model=tl_model, - output_type=output_type, - verbose=verbose, - is_classifier=is_classifier, - default_chunk_size=default_chunk_size, - align_bytes=align_bytes, - layout=layout, - precision=precision, - device_id=device_id - ) - - @nvtx.annotate( - message='ForestInference.predict_proba', - domain='cuml_python' - ) - @reflect - def predict_proba( - self, - X, - *, - preds=None, - chunk_size=None, - ) -> CumlArray: - """ - Predict the class probabilities for each row in X. - - Parameters - ---------- - X - The input data of shape Rows X Features. This can be a numpy - array, cupy array, Pandas/cuDF Dataframe or any other array type - accepted by cuML. FIL is optimized for C-major arrays (e.g. - numpy/cupy arrays). Inputs whose datatype does not match the - precision of the loaded model (float/double) will be converted - to the correct datatype before inference. If this input is in a - memory location that is inaccessible to the current device type - (as set with e.g. the `set_fil_device_type` context manager), - it will be copied to the correct location. This copy will be - distributed across as many CUDA streams as are available - in the stream pool of the model's RAFT handle. - preds - If non-None, outputs will be written in-place to this array. - Therefore, if given, this should be a C-major array of shape Rows x - Classes with a datatype (float/double) corresponding to the - precision of the model. If None, an output array of the correct - shape and type will be allocated and returned. - chunk_size : int - The number of rows to simultaneously process in one iteration - of the inference algorithm. Batches are further broken down into - "chunks" of this size when assigning available threads to tasks. - The choice of chunk size can have a substantial impact on - performance, but the optimal choice depends on model and - hardware and is difficult to predict a priori. In general, - larger batch sizes benefit from larger chunk sizes, and smaller - batch sizes benefit from small chunk sizes. On GPU, valid - values are powers of 2 from 1 to 32. On CPU, valid values are - any power of 2, but little benefit is expected above a chunk size - of 512. - """ - if not self.is_classifier: - raise RuntimeError( - "predict_proba is not available for regression models. Load" - " with is_classifier=True if this is a classifier." - ) - return self.forest.predict( - X, preds=preds, chunk_size=(chunk_size or self.default_chunk_size) - ) - - @nvtx.annotate( - message='ForestInference.predict', - domain='cuml_python' - ) - @reflect - def predict( - self, - X, - *, - preds=None, - chunk_size=None, - threshold=None, - ) -> CumlArray: - """ - For classification models, predict the class for each row. For - regression models, predict the output for each row. - - Parameters - ---------- - X - The input data of shape Rows X Features. This can be a numpy - array, cupy array, Pandas/cuDF Dataframe or any other array type - accepted by cuML. FIL is optimized for C-major arrays (e.g. - numpy/cupy arrays). Inputs whose datatype does not match the - precision of the loaded model (float/double) will be converted - to the correct datatype before inference. If this input is in a - memory location that is inaccessible to the current device type - (as set with e.g. the `set_fil_device_type` context manager), - it will be copied to the correct location. This copy will be - distributed across as many CUDA streams as are available - in the stream pool of the model's RAFT handle. - preds - If non-None, outputs will be written in-place to this array. - Therefore, if given, this should be a C-major array of shape Rows x - 1 with a datatype (float/double) corresponding to the precision of - the model. If None, an output array of the correct shape and - type will be allocated and returned. For classifiers, in-place - prediction offers no performance or memory benefit. For regressors, - in-place prediction offers both a performance and memory - benefit. - chunk_size : int - The number of rows to simultaneously process in one iteration - of the inference algorithm. Batches are further broken down into - "chunks" of this size when assigning available threads to tasks. - The choice of chunk size can have a substantial impact on - performance, but the optimal choice depends on model and - hardware and is difficult to predict a priori. In general, - larger batch sizes benefit from larger chunk sizes, and smaller - batch sizes benefit from small chunk sizes. On GPU, valid - values are powers of 2 from 1 to 32. On CPU, valid values are - any power of 2, but little benefit is expected above a chunk size - of 512. - threshold : float - For binary classifiers, output probabilities above this threshold - will be considered positive detections. If None, a threshold - of 0.5 will be used for binary classifiers. For multiclass - classifiers, the highest probability class is chosen regardless - of threshold. - """ - chunk_size = (chunk_size or self.default_chunk_size) - if self.forest.row_postprocessing() == 'max_index': - raw_out = self.forest.predict(X, chunk_size=chunk_size) - result = raw_out[:, 0] - if preds is None: - return result - else: - preds[:] = result - return preds - elif self.is_classifier: - proba = self.forest.predict(X, chunk_size=chunk_size) - if len(proba.shape) < 2 or proba.shape[1] == 1: - if threshold is None: - threshold = 0.5 - result = ( - proba.to_output(output_type='array') > threshold - ).astype('int') - else: - result = GlobalSettings().fil_xpy.argmax( - proba.to_output(output_type='array'), axis=1 - ) - if preds is None: - return CumlArray(data=result, index=proba.index) - else: - preds[:] = result - return preds - else: - return self.forest.predict( - X, predict_type="default", preds=preds, chunk_size=chunk_size - ) - - @nvtx.annotate( - message='ForestInference.predict_per_tree', - domain='cuml_python' - ) - @reflect - def predict_per_tree( - self, - X, - *, - preds=None, - chunk_size=None) -> CumlArray: - """ - Output prediction of each tree. - This function computes one or more margin scores per tree. - - Parameters - ---------- - X - The input data of shape Rows X Features. This can be a numpy - array, cupy array, Pandas/cuDF Dataframe or any other array type - accepted by cuML. FIL is optimized for C-major arrays (e.g. - numpy/cupy arrays). Inputs whose datatype does not match the - precision of the loaded model (float/double) will be converted - to the correct datatype before inference. If this input is in a - memory location that is inaccessible to the current device type - (as set with e.g. the `set_fil_device_type` context manager), - it will be copied to the correct location. This copy will be - distributed across as many CUDA streams as are available - in the stream pool of the model's RAFT handle. - preds - If non-None, outputs will be written in-place to this array. - Therefore, if given, this should be a C-major array of shape - n_rows * n_trees * n_outputs (if vector leaf is used) or - shape n_rows * n_trees (if scalar leaf is used). - Classes with a datatype (float/double) corresponding to the - precision of the model. If None, an output array of the correct - shape and type will be allocated and returned. - chunk_size : int - The number of rows to simultaneously process in one iteration - of the inference algorithm. Batches are further broken down into - "chunks" of this size when assigning available threads to tasks. - The choice of chunk size can have a substantial impact on - performance, but the optimal choice depends on model and - hardware and is difficult to predict a priori. In general, - larger batch sizes benefit from larger chunk sizes, and smaller - batch sizes benefit from small chunk sizes. On GPU, valid - values are powers of 2 from 1 to 32. On CPU, valid values are - any power of 2, but little benefit is expected above a chunk size - of 512. - """ - chunk_size = (chunk_size or self.default_chunk_size) - return self.forest.predict( - X, predict_type="per_tree", preds=preds, chunk_size=chunk_size - ) - - @nvtx.annotate( - message='ForestInference.apply', - domain='cuml_python' - ) - @reflect - def apply( - self, - X, - *, - preds=None, - chunk_size=None) -> CumlArray: - """ - Output the ID of the leaf node for each tree. - - Parameters - ---------- - X - The input data of shape Rows X Features. This can be a numpy - array, cupy array, Pandas/cuDF Dataframe or any other array type - accepted by cuML. FIL is optimized for C-major arrays (e.g. - numpy/cupy arrays). Inputs whose datatype does not match the - precision of the loaded model (float/double) will be converted - to the correct datatype before inference. If this input is in a - memory location that is inaccessible to the current device type - (as set with e.g. the `set_fil_device_type` context manager), - it will be copied to the correct location. This copy will be - distributed across as many CUDA streams as are available - in the stream pool of the model's RAFT handle. - preds - If non-None, outputs will be written in-place to this array. - Therefore, if given, this should be a C-major array of shape - n_rows * n_trees. - Classes with a datatype (float/double) corresponding to the - precision of the model. If None, an output array of the correct - shape and type will be allocated and returned. - chunk_size : int - The number of rows to simultaneously process in one iteration - of the inference algorithm. Batches are further broken down into - "chunks" of this size when assigning available threads to tasks. - The choice of chunk size can have a substantial impact on - performance, but the optimal choice depends on model and - hardware and is difficult to predict a priori. In general, - larger batch sizes benefit from larger chunk sizes, and smaller - batch sizes benefit from small chunk sizes. On GPU, valid - values are powers of 2 from 1 to 32. On CPU, valid values are - any power of 2, but little benefit is expected above a chunk size - of 512. - """ - return self.forest.predict( - X, predict_type="leaf_id", preds=preds, chunk_size=chunk_size - ) - - def optimize( - self, - *, - data=None, - batch_size=1024, - unique_batches=10, - timeout=0.2, - predict_method='predict', - max_chunk_size=None, - seed=0 - ): - """ - Find the optimal layout and chunk size for this model - - The optimal value for layout and chunk size depends on the model, - batch size, and available hardware. In order to get the most - realistic performance distribution, example data can be provided. If - it is not, random data will be generated based on the indicated batch - size. After finding the optimal layout, the model will be reloaded if - necessary. The optimal chunk size will be used to set the default chunk - size used if none is passed to the predict call. - - Parameters - ---------- - data - Example data either of shape unique_batches x batch size x features - or batch_size x features or None. If None, random data will be - generated instead. - batch_size : int - If example data is not provided, random data with this many rows - per batch will be used. - unique_batches : int - The number of unique batches to generate if random data are used. - Increasing this number decreases the chance that the optimal - configuration will be skewed by a single batch with unusual - performance characteristics. - timeout : float - Time in seconds to target for optimization. The optimization loop - will be repeatedly run a number of times increasing in the sequence - 1, 2, 5, 10, 20, 50, ... until the time taken is at least the given - value. Note that for very large batch sizes and large models, the - total elapsed time may exceed this timeout; it is a soft target for - elapsed time. Setting the timeout to zero will run through the - indicated number of unique batches exactly once. Defaults to 0.2s. - predict_method : str - If desired, optimization can occur over one of the prediction - method variants (e.g. "predict_per_tree") rather than the - default `predict` method. To do so, pass the name of the method - here. - max_chunk_size : int or None - The maximum chunk size to explore during optimization. If not - set, a value will be picked based on the current device type. - Setting this to a lower value will reduce the optimization search - time but may not result in optimal performance. - seed : int - The random seed used for generating example data if none is - provided. - """ - if data is None: - xpy = GlobalSettings().fil_xpy - dtype = self.forest.get_dtype() - data = xpy.random.uniform( - xpy.finfo(dtype).min / 2, - xpy.finfo(dtype).max / 2, - (unique_batches, batch_size, self.forest.num_features()) - ) - else: - data = CumlArray.from_input( - data, - order='K', - convert_to_mem_type=GlobalSettings().fil_memory_type, - ).to_output('array') - try: - unique_batches, batch_size, features = data.shape - except ValueError: - unique_batches = 1 - batch_size, features = data.shape - data = [data] - - if max_chunk_size is None: - max_chunk_size = 512 - if get_fil_device_type() is DeviceType.device: - max_chunk_size = min(max_chunk_size, 32) - - max_chunk_size = min(max_chunk_size, batch_size) - - infer = getattr(self, predict_method) - - optimal_layout = 'depth_first' - optimal_chunk_size = 1 - - valid_layouts = ('depth_first', 'breadth_first', 'layered') - chunk_size = 1 - valid_chunk_sizes = [] - while chunk_size <= max_chunk_size: - valid_chunk_sizes.append(chunk_size) - chunk_size *= 2 - - all_params = list(itertools.product(valid_layouts, valid_chunk_sizes)) - auto_iterator = _AutoIterations() - loop_start = perf_counter() - while True: - optimal_time = float('inf') - iterations = auto_iterator.next() - for layout, chunk_size in all_params: - self.layout = layout - infer(data[0], chunk_size=chunk_size) - elapsed = float('inf') - for _ in range(iterations): - start = perf_counter() - for iter_index in range(unique_batches): - infer( - data[iter_index], chunk_size=chunk_size - ) - elapsed = min(elapsed, perf_counter() - start) - if elapsed < optimal_time: - optimal_time = elapsed - optimal_layout = layout - optimal_chunk_size = chunk_size - if (perf_counter() - loop_start > timeout): - break - - self.layout = optimal_layout - self.default_chunk_size = optimal_chunk_size - - @classmethod - def _get_param_names(cls): - return [ - *super()._get_param_names(), - "treelite_model", - "is_classifier", - "layout", - "default_chunk_size", - "align_bytes", - "precision", - "device_id", - "ensure_all_finite", - ] diff --git a/python/cuml/cuml/fil/infer_kind.pxd b/python/cuml/cuml/fil/infer_kind.pxd deleted file mode 100644 index c3bc1bd9fd..0000000000 --- a/python/cuml/cuml/fil/infer_kind.pxd +++ /dev/null @@ -1,11 +0,0 @@ -# -# SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# - -cdef extern from "cuml/fil/infer_kind.hpp" namespace "ML::fil": - # TODO(hcho3): Switch to new syntax for scoped enum when we adopt Cython 3.0 - cdef enum infer_kind: - default_kind "ML::fil::infer_kind::default_kind" - per_tree "ML::fil::infer_kind::per_tree" - leaf_id "ML::fil::infer_kind::leaf_id" diff --git a/python/cuml/cuml/fil/postprocessing.pxd b/python/cuml/cuml/fil/postprocessing.pxd deleted file mode 100644 index da1f08f762..0000000000 --- a/python/cuml/cuml/fil/postprocessing.pxd +++ /dev/null @@ -1,16 +0,0 @@ -# -# SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# -cdef extern from "cuml/fil/postproc_ops.hpp" namespace "ML::fil" nogil: - cdef enum row_op: - row_disable "ML::fil::row_op::disable", - softmax "ML::fil::row_op::softmax", - max_index "ML::fil::row_op::max_index" - cdef enum element_op: - elem_disable "ML::fil::element_op::disable", - signed_square "ML::fil::element_op::signed_square", - hinge "ML::fil::element_op::hinge", - sigmoid "ML::fil::element_op::sigmoid", - exponential "ML::fil::element_op::exponential", - logarithm_one_plus_exp "ML::fil::element_op::logarithm_one_plus_exp" diff --git a/python/cuml/cuml/fil/tree_layout.pxd b/python/cuml/cuml/fil/tree_layout.pxd deleted file mode 100644 index 0b73171a5c..0000000000 --- a/python/cuml/cuml/fil/tree_layout.pxd +++ /dev/null @@ -1,9 +0,0 @@ -# -# SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. -# SPDX-License-Identifier: Apache-2.0 -# -cdef extern from "cuml/fil/tree_layout.hpp" namespace "ML::fil" nogil: - cdef enum tree_layout: - depth_first "ML::fil::tree_layout::depth_first", - breadth_first "ML::fil::tree_layout::breadth_first", - layered_children_together "ML::fil::tree_layout::layered_children_together" diff --git a/python/cuml/pyproject.toml b/python/cuml/pyproject.toml index 04ae556bb7..001dd390b0 100644 --- a/python/cuml/pyproject.toml +++ b/python/cuml/pyproject.toml @@ -88,6 +88,7 @@ dependencies = [ "numba-cuda>=0.22.2,<0.29.0", "numba>=0.60.0,<0.65.0", "numpy>=1.26,<3.0", + "nvforest==26.6.*,>=0.0.0a0", "nvidia-nvjitlink>=13.0,<14", "packaging", "pylibraft==26.6.*,>=0.0.0a0", @@ -186,9 +187,11 @@ requires = [ "cuda-python>=13.0.1,<14.0", "cython>=3.2.2", "libcuml==26.6.*,>=0.0.0a0", + "libnvforest==26.6.*,>=0.0.0a0", "libraft==26.6.*,>=0.0.0a0", "librmm==26.6.*,>=0.0.0a0", "ninja", + "nvforest==26.6.*,>=0.0.0a0", "pylibraft==26.6.*,>=0.0.0a0", "rmm==26.6.*,>=0.0.0a0", "treelite>=4.7.0,<5.0.0", diff --git a/python/cuml/tests/explainer/test_explainer_common.py b/python/cuml/tests/explainer/test_explainer_common.py index a5822270dd..c752fd8a3b 100644 --- a/python/cuml/tests/explainer/test_explainer_common.py +++ b/python/cuml/tests/explainer/test_explainer_common.py @@ -21,10 +21,16 @@ ) from cuml.testing.utils import ClassEnumerator -# TODO(26.08) Remove this filter -pytestmark = pytest.mark.filterwarnings( - "ignore:The default value of 'max_depth':FutureWarning" -) +pytestmark = [ + # TODO(26.10) Remove this filter, once cuml.fil is removed + pytest.mark.filterwarnings( + "ignore:cuml.fil.ForestInference.* is deprecated:FutureWarning" + ), + # TODO(26.08) Remove this filter + pytest.mark.filterwarnings( + "ignore:The default value of 'max_depth':FutureWarning" + ), +] models_config = ClassEnumerator(module=cuml) models = models_config.get_models() diff --git a/python/cuml/tests/test_api.py b/python/cuml/tests/test_api.py index 4a3c4854aa..f1aeeca297 100644 --- a/python/cuml/tests/test_api.py +++ b/python/cuml/tests/test_api.py @@ -15,6 +15,11 @@ from cuml.internals.base import Base from cuml.testing.utils import ClassEnumerator +# TODO(26.10) Remove this filter, once cuml.fil is removed +pytestmark = pytest.mark.filterwarnings( + "ignore:cuml.fil.ForestInference.* is deprecated:FutureWarning" +) + ############################################################################### # Helper functions and classes # ############################################################################### diff --git a/python/cuml/tests/test_benchmark.py b/python/cuml/tests/test_benchmark.py index 17da5a3193..d35f358809 100644 --- a/python/cuml/tests/test_benchmark.py +++ b/python/cuml/tests/test_benchmark.py @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # import time @@ -176,7 +176,6 @@ def predict(self, X): "DBSCAN", "LogisticRegression", "ElasticNet", - "FIL", "xgboost-classification", "xgboost-regression", ], @@ -184,7 +183,7 @@ def predict(self, X): def test_real_algos_runner(algo_name): pair = algorithms.algorithm_by_name(algo_name) - if algo_name in ["FIL", "xgboost-classification", "xgboost-regression"]: + if algo_name in ["xgboost-classification", "xgboost-regression"]: pytest.importorskip("xgboost") # Use appropriate dataset for regression algorithms @@ -201,26 +200,6 @@ def test_real_algos_runner(algo_name): assert results["cuml_acc"] is not None -# Test FIL with several input types -@pytest.mark.parametrize( - "input_type", ["numpy", "cudf", "gpuarray", "gpuarray-c"] -) -def test_fil_input_types(input_type): - pair = algorithms.algorithm_by_name("FIL") - - pytest.importorskip("xgboost") - - runner = AccuracyComparisonRunner( - [20], - [5], - dataset_name="classification", - test_fraction=0.5, - input_type=input_type, - ) - results = runner.run(pair, run_cpu=False)[0] - assert results["cuml_acc"] is not None - - @pytest.mark.parametrize("input_type", ["numpy", "cudf", "pandas", "gpuarray"]) def test_training_data_to_numpy(input_type): X, y, *_ = datagen.gen_data( diff --git a/python/cuml/tests/test_fil.py b/python/cuml/tests/test_fil.py index da1df49a74..7cdb281bd9 100644 --- a/python/cuml/tests/test_fil.py +++ b/python/cuml/tests/test_fil.py @@ -10,9 +10,9 @@ import numpy as np import pandas as pd import pytest -import sklearn import treelite -from packaging.version import Version + +# TODO(26.10): Remove this once we fully phase out cuml.fil.ForestInference # Import XGBoost before scikit-learn to work around a libgomp bug # See https://github.com/dmlc/xgboost/issues/7110 @@ -42,10 +42,20 @@ unit_param, ) -# TODO(26.08): Remove this filter -pytestmark = pytest.mark.filterwarnings( - "ignore:The default value of 'max_depth':FutureWarning" -) +pytestmark = [ + pytest.mark.filterwarnings( + "ignore:cuml.fil.ForestInference.* is deprecated:FutureWarning" + ), + pytest.mark.filterwarnings( + "ignore:.*set_fil_device_type is deprecated:FutureWarning" + ), + pytest.mark.filterwarnings( + r"ignore:.*as_fil\(\) method is deprecated.*:FutureWarning" + ), + pytest.mark.filterwarnings( + "ignore:The default value of 'max_depth':FutureWarning" + ), +] def simulate_data( @@ -679,7 +689,7 @@ def test_lightgbm( @pytest.mark.parametrize("train_device", ("cpu", "gpu")) @pytest.mark.parametrize("infer_device", ("cpu", "gpu")) @pytest.mark.parametrize("n_classes", [2, 5, 25]) -@pytest.mark.parametrize("num_boost_round", [10, 100]) +@pytest.mark.parametrize("num_boost_round", [10, 20]) def test_predict_per_tree( train_device, infer_device, n_classes, num_boost_round, tmp_path ): @@ -834,7 +844,7 @@ def test_missing_categorical(category_list): leaf_output_type="float32", metadata=treelite.model_builder.Metadata( num_feature=1, - task_type="kBinaryClf", + task_type="kRegressor", average_tree_output=False, num_target=1, num_class=[1], @@ -917,9 +927,6 @@ def test_device_selection(device_id, model_kind, tmp_path): ) xgb_model.fit(X, y) model_path = os.path.join(tmp_path, "xgb_class.ubj") - # skip with sklearn version 1.8.0.dev0 - if Version(sklearn.__version__) >= Version("1.8.0.dev0"): - pytest.skip("xgboost is incompatible with sklearn >= 1.8.0.dev0") xgb_model.save_model(model_path) fm = ForestInference.load( model_path, @@ -981,7 +988,7 @@ def test_wide_data(): n_rows = 50 n_features = 100000 X = np.random.normal(size=(n_rows, n_features)).astype(np.float32) - y = np.asarray([0, 1] * (n_rows // 2), dtype=np.int32) + y = np.array([0, 1] * (n_rows // 2), dtype=np.int32) clf = RandomForestClassifier(max_features="sqrt", n_estimators=10) clf.fit(X, y) diff --git a/python/cuml/tests/test_random_forest.py b/python/cuml/tests/test_random_forest.py index 5c1e9c043e..efaf3d01ac 100644 --- a/python/cuml/tests/test_random_forest.py +++ b/python/cuml/tests/test_random_forest.py @@ -9,6 +9,7 @@ import warnings import cudf +import cupy as cp import numpy as np import pytest import treelite @@ -507,6 +508,14 @@ def test_rf_classification_fit_and_predict_dtypes_differ( cuml_model = curfc() cuml_model.fit(X_train, y_train) + + if not convert_dtype: + with pytest.raises( + ValueError, match=r".*Expected array with dtype in.*" + ): + preds = cuml_model.predict(X_test, convert_dtype=convert_dtype) + return + preds = cuml_model.predict(X_test, convert_dtype=convert_dtype) acc = accuracy_score(y_test, preds) if X.shape[0] < 500000: @@ -647,14 +656,14 @@ def test_rf_classification_proba( @pytest.mark.parametrize("datatype", [np.float32, np.float64]) @pytest.mark.parametrize( - "fil_layout", ["depth_first", "breadth_first", "layered"] + "nvforest_layout", ["depth_first", "breadth_first", "layered"] ) @pytest.mark.skipif( cudf_pandas_active, reason="cudf.pandas causes sklearn RF estimators crashes sometimes. " "Issue: https://github.com/rapidsai/cuml/issues/5991", ) -def test_rf_classification_sparse(small_clf, datatype, fil_layout): +def test_rf_classification_sparse(small_clf, datatype, nvforest_layout): num_trees = 50 X, y = small_clf @@ -677,16 +686,16 @@ def test_rf_classification_sparse(small_clf, datatype, fil_layout): max_depth=40, ) cuml_model.fit(X_train, y_train) - preds = cuml_model.predict(X_test, layout=fil_layout) + preds = cuml_model.predict(X_test, layout=nvforest_layout) acc = accuracy_score(y_test, preds) np.testing.assert_almost_equal(acc, cuml_model.score(X_test, y_test)) - fil_model = cuml_model.as_fil() + nvforest_model = cuml_model.as_nvforest(layout=nvforest_layout) with cuml.using_output_type("numpy"): - fil_model_preds = fil_model.predict(X_test) - fil_model_acc = accuracy_score(y_test, fil_model_preds) - assert acc == fil_model_acc + nvforest_model_preds = cp.asnumpy(nvforest_model.predict(X_test)) + nvforest_model_acc = accuracy_score(y_test, nvforest_model_preds) + assert acc == nvforest_model_acc tl_model = cuml_model.as_treelite() assert num_trees == tl_model.num_tree @@ -709,14 +718,14 @@ def test_rf_classification_sparse(small_clf, datatype, fil_layout): @pytest.mark.parametrize("datatype", [np.float32, np.float64]) @pytest.mark.parametrize( - "fil_layout", ["depth_first", "breadth_first", "layered"] + "nvforest_layout", ["depth_first", "breadth_first", "layered"] ) @pytest.mark.skipif( cudf_pandas_active, reason="cudf.pandas causes sklearn RF estimators crashes sometimes. " "Issue: https://github.com/rapidsai/cuml/issues/5991", ) -def test_rf_regression_sparse(special_reg, datatype, fil_layout): +def test_rf_regression_sparse(special_reg, datatype, nvforest_layout): num_trees = 50 X, y = special_reg @@ -739,16 +748,16 @@ def test_rf_regression_sparse(special_reg, datatype, fil_layout): ) cuml_model.fit(X_train, y_train) - preds = cuml_model.predict(X_test, layout=fil_layout) + preds = cuml_model.predict(X_test, layout=nvforest_layout) r2 = r2_score(y_test, preds) - fil_model = cuml_model.as_fil() + nvforest_model = cuml_model.as_nvforest(layout=nvforest_layout) - with cuml.using_output_type("numpy"): - fil_model_preds = fil_model.predict(X_test) - fil_model_preds = np.reshape(fil_model_preds, np.shape(y_test)) - fil_model_r2 = r2_score(y_test, fil_model_preds) - assert r2 == fil_model_r2 + nvforest_model_preds = cp.reshape( + nvforest_model.predict(X_test), np.shape(y_test) + ) + nvforest_model_r2 = r2_score(cp.asarray(y_test), nvforest_model_preds) + assert r2 == nvforest_model_r2 tl_model = cuml_model.as_treelite() assert num_trees == tl_model.num_tree diff --git a/python/libcuml/libcuml/load.py b/python/libcuml/libcuml/load.py index 5aa38f02d2..8ab00af389 100644 --- a/python/libcuml/libcuml/load.py +++ b/python/libcuml/libcuml/load.py @@ -37,11 +37,13 @@ def load_library(): """Dynamically load libcuml.so and its dependencies""" try: # These libraries must all be loaded before libcuml + import libnvforest import libraft import librmm import rapids_logger rapids_logger.load_library() + libnvforest.load_library() librmm.load_library() libraft.load_library() except ModuleNotFoundError: diff --git a/python/libcuml/pyproject.toml b/python/libcuml/pyproject.toml index 523df2d06f..8dbf01d398 100644 --- a/python/libcuml/pyproject.toml +++ b/python/libcuml/pyproject.toml @@ -26,6 +26,7 @@ classifiers = [ ] dependencies = [ "cuda-toolkit[cublas,cufft,curand,cusolver,cusparse]==13.*", + "libnvforest==26.6.*,>=0.0.0a0", "libraft==26.6.*,>=0.0.0a0", "librmm==26.6.*,>=0.0.0a0", "nvidia-nvjitlink>=13.0,<14", @@ -67,6 +68,7 @@ dependencies-file = "../../dependencies.yaml" matrix-entry = "cuda_suffixed=true;use_cuda_wheels=true" requires = [ "cmake>=4.0", + "libnvforest==26.6.*,>=0.0.0a0", "libraft==26.6.*,>=0.0.0a0", "librmm==26.6.*,>=0.0.0a0", "ninja", From 65c3d1e838337df30951843faf33acb96d39ce2e Mon Sep 17 00:00:00 2001 From: Simon Adorf Date: Thu, 21 May 2026 21:59:37 +0000 Subject: [PATCH 17/17] Suppress intentional hardcoded version references --- python/cuml/cuml/svm/linear_svc.py | 1 + python/cuml/cuml/svm/svc.py | 1 + python/cuml/cuml_accel_tests/integration/test_svc.py | 1 + python/cuml/cuml_accel_tests/upstream/pytest.ini | 1 + python/cuml/tests/explainer/test_explainer_permutation_shap.py | 1 + python/cuml/tests/test_linear_svm.py | 1 + python/cuml/tests/test_sklearn_import_export.py | 1 + python/cuml/tests/test_svm.py | 1 + 8 files changed, 8 insertions(+) diff --git a/python/cuml/cuml/svm/linear_svc.py b/python/cuml/cuml/svm/linear_svc.py index c13f30e9ae..80b961cd9e 100644 --- a/python/cuml/cuml/svm/linear_svc.py +++ b/python/cuml/cuml/svm/linear_svc.py @@ -255,6 +255,7 @@ def fit( if self.probability != "deprecated": warnings.warn( "The `probability` parameter is deprecated and will be " + # rapids-pre-commit-hooks: disable-next-line "removed in cuML version 26.08. Use " "`CalibratedClassifierCV(LinearSVC(), ensemble=False)` from " "`sklearn.calibration` instead.", diff --git a/python/cuml/cuml/svm/svc.py b/python/cuml/cuml/svm/svc.py index bd5295b0ec..778a3f328d 100644 --- a/python/cuml/cuml/svm/svc.py +++ b/python/cuml/cuml/svm/svc.py @@ -459,6 +459,7 @@ def fit(self, X, y, sample_weight=None, *, convert_dtype=True) -> "SVC": if self.probability != "deprecated": warnings.warn( "The `probability` parameter is deprecated and will be " + # rapids-pre-commit-hooks: disable-next-line "removed in cuML version 26.08. Use " "`CalibratedClassifierCV(SVC(), ensemble=False)` from " "`sklearn.calibration` instead.", diff --git a/python/cuml/cuml_accel_tests/integration/test_svc.py b/python/cuml/cuml_accel_tests/integration/test_svc.py index 642e3463a0..52411e2e8b 100644 --- a/python/cuml/cuml_accel_tests/integration/test_svc.py +++ b/python/cuml/cuml_accel_tests/integration/test_svc.py @@ -37,6 +37,7 @@ def test_svc(binary): assert svc.score(X, y) > 0.5 +# rapids-pre-commit-hooks: disable-next-line # TODO(26.08): Remove once `probability` is removed from cuml.svm.SVC. @pytest.mark.filterwarnings( "ignore:Attribute `prob[AB]_` was deprecated:FutureWarning" diff --git a/python/cuml/cuml_accel_tests/upstream/pytest.ini b/python/cuml/cuml_accel_tests/upstream/pytest.ini index c895ea73f2..8c406ae261 100644 --- a/python/cuml/cuml_accel_tests/upstream/pytest.ini +++ b/python/cuml/cuml_accel_tests/upstream/pytest.ini @@ -16,6 +16,7 @@ filterwarnings = error::FutureWarning:cuml # Suppress cuml's deprecation warning for sklearn upstream tests that # legitimately use `probability=True`. Must come after the generic error + # rapids-pre-commit-hooks: disable-next-line # rule (later filterwarnings entries take precedence). TODO(26.08): drop. ignore:The `probability` parameter is deprecated:FutureWarning # Ignore unknown pytest marks, the xfail-list currently adds a bunch of these diff --git a/python/cuml/tests/explainer/test_explainer_permutation_shap.py b/python/cuml/tests/explainer/test_explainer_permutation_shap.py index b80b44b24a..9783adbb71 100644 --- a/python/cuml/tests/explainer/test_explainer_permutation_shap.py +++ b/python/cuml/tests/explainer/test_explainer_permutation_shap.py @@ -54,6 +54,7 @@ def test_regression_datasets(exact_shap_regression_dataset, model): ) <= 1e-5 +# rapids-pre-commit-hooks: disable-next-line # TODO(26.08): Remove this filter once `probability` is removed from cuml.svm.SVC. @pytest.mark.filterwarnings( "ignore:The `probability` parameter is deprecated:FutureWarning" diff --git a/python/cuml/tests/test_linear_svm.py b/python/cuml/tests/test_linear_svm.py index 41182aecb4..bdaa475831 100644 --- a/python/cuml/tests/test_linear_svm.py +++ b/python/cuml/tests/test_linear_svm.py @@ -214,6 +214,7 @@ def test_linear_svc_decision_function( @pytest.mark.parametrize("fit_intercept", [True, False]) @pytest.mark.parametrize("n_classes", [2, 3, 5]) +# rapids-pre-commit-hooks: disable-next-line # TODO(26.08): Remove once `probability` is removed from cuml.svm.LinearSVC. @pytest.mark.filterwarnings( "ignore:The `probability` parameter is deprecated:FutureWarning" diff --git a/python/cuml/tests/test_sklearn_import_export.py b/python/cuml/tests/test_sklearn_import_export.py index 92f2aa780d..704585c432 100644 --- a/python/cuml/tests/test_sklearn_import_export.py +++ b/python/cuml/tests/test_sklearn_import_export.py @@ -376,6 +376,7 @@ def test_svr(random_state, sparse, kernel): ) +# rapids-pre-commit-hooks: disable-next-line # TODO(26.08): Remove this filter once `probability` is removed from cuml.svm.SVC. @pytest.mark.filterwarnings( "ignore:The `probability` parameter (is|was) deprecated:FutureWarning" diff --git a/python/cuml/tests/test_svm.py b/python/cuml/tests/test_svm.py index 107c7e622d..0a2ee6bb87 100644 --- a/python/cuml/tests/test_svm.py +++ b/python/cuml/tests/test_svm.py @@ -40,6 +40,7 @@ # Many tests below pass `probability=` to cuml SVC/LinearSVC on purpose; # silence the FutureWarning module-wide. +# rapids-pre-commit-hooks: disable-next-line # TODO(26.08): Remove once `probability` is removed from cuml.svm.SVC/LinearSVC. pytestmark = pytest.mark.filterwarnings( "ignore:The `probability` parameter is deprecated:FutureWarning"