Skip to content

Commit b9030fc

Browse files
authored
Merge branch 'main' into fea-bench-yaml
2 parents ea3c199 + 146f616 commit b9030fc

17 files changed

Lines changed: 2438 additions & 228 deletions

File tree

cpp/examples/symreg/symreg_example.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -235,7 +235,8 @@ int main(int argc, char* argv[])
235235

236236
// Initialize AST
237237
auto curr_mr = rmm::mr::get_current_device_resource_ref();
238-
d_finalprogs = static_cast<cg::program_t>(curr_mr.allocate(stream, params.population_size));
238+
d_finalprogs = static_cast<cg::program_t>(
239+
curr_mr.allocate(stream, params.population_size * sizeof(cg::program), alignof(cg::program)));
239240

240241
std::vector<std::vector<cg::program>> history;
241242
history.reserve(params.generations);
@@ -327,7 +328,8 @@ int main(int argc, char* argv[])
327328

328329
/* ======================= Reset data ======================= */
329330

330-
curr_mr.deallocate(stream, d_finalprogs, params.population_size);
331+
curr_mr.deallocate(
332+
stream, d_finalprogs, params.population_size * sizeof(cg::program), alignof(cg::program));
331333
CUDA_RT_CALL(cudaEventDestroy(start));
332334
CUDA_RT_CALL(cudaEventDestroy(stop));
333335
return 0;

cpp/src/genetic/genetic.cu

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -221,15 +221,15 @@ void parallel_evolve(const raft::handle_t& h,
221221

222222
// Set current generation device nodes
223223
tmp.nodes = (node*)rmm::mr::get_current_device_resource_ref().allocate(
224-
stream, h_nextprogs[i].len * sizeof(node));
224+
stream, h_nextprogs[i].len * sizeof(node), alignof(node));
225225
raft::copy(tmp.nodes, h_nextprogs[i].nodes, h_nextprogs[i].len, stream);
226226
raft::copy(d_nextprogs + i, &tmp, 1, stream);
227227

228228
if (generation > 1) {
229229
// Free device memory allocated to program nodes in previous generation
230230
raft::copy(&tmp, d_oldprogs + i, 1, stream);
231231
rmm::mr::get_current_device_resource_ref().deallocate(
232-
stream, tmp.nodes, h_nextprogs[i].len * sizeof(node));
232+
stream, tmp.nodes, tmp.len * sizeof(node), alignof(node));
233233
}
234234

235235
tmp.nodes = nullptr;
@@ -399,7 +399,7 @@ void symFit(const raft::handle_t& handle,
399399

400400
program_t d_currprogs; // pointer to current programs
401401
d_currprogs = (program_t)rmm::mr::get_current_device_resource_ref().allocate(
402-
stream, params.population_size * sizeof(program));
402+
stream, params.population_size * sizeof(program), alignof(program));
403403
program_t d_nextprogs = final_progs; // Reuse memory already allocated for final_progs
404404
final_progs = nullptr;
405405

@@ -481,7 +481,7 @@ void symFit(const raft::handle_t& handle,
481481

482482
// Deallocate the previous generation device memory
483483
rmm::mr::get_current_device_resource_ref().deallocate(
484-
stream, d_nextprogs, params.population_size * sizeof(program));
484+
stream, d_nextprogs, params.population_size * sizeof(program), alignof(program));
485485
d_currprogs = nullptr;
486486
d_nextprogs = nullptr;
487487
}

cpp/tests/sg/genetic/evolution_test.cu

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -259,7 +259,7 @@ TEST_F(GeneticEvolutionTest, SymReg)
259259
MLCommon::CompareApprox<float> compApprox(tolerance);
260260
program_t final_progs;
261261
final_progs = (program_t)rmm::mr::get_current_device_resource_ref().allocate(
262-
stream, hyper_params.population_size * sizeof(program));
262+
stream, hyper_params.population_size * sizeof(program), alignof(program));
263263
std::vector<std::vector<program>> history;
264264
history.reserve(hyper_params.generations);
265265

@@ -327,12 +327,12 @@ TEST_F(GeneticEvolutionTest, SymReg)
327327
program tmp = program();
328328
raft::copy(&tmp, final_progs + i, 1, stream);
329329
rmm::mr::get_current_device_resource_ref().deallocate(
330-
stream, tmp.nodes, tmp.len * sizeof(node));
330+
stream, tmp.nodes, tmp.len * sizeof(node), alignof(node));
331331
tmp.nodes = nullptr;
332332
}
333333
// deallocate the final programs from device memory
334334
rmm::mr::get_current_device_resource_ref().deallocate(
335-
stream, final_progs, hyper_params.population_size * sizeof(program));
335+
stream, final_progs, hyper_params.population_size * sizeof(program), alignof(program));
336336

337337
ASSERT_TRUE(compApprox(history[n_gen - 1][best_idx].raw_fitness_, 0.0036f));
338338
std::cout << "Some Predicted test values:" << std::endl;

cpp/tests/sg/genetic/program_test.cu

Lines changed: 12 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -90,10 +90,12 @@ class GeneticProgramTest : public ::testing::Test {
9090
d_lY.resize(250, stream);
9191
d_lunitW.resize(250, stream);
9292
d_lW.resize(250, stream);
93-
d_nodes1 = (node*)rmm::mr::get_current_device_resource_ref().allocate(stream, 7 * sizeof(node));
94-
d_nodes2 = (node*)rmm::mr::get_current_device_resource_ref().allocate(stream, 7 * sizeof(node));
95-
d_progs =
96-
(program_t)rmm::mr::get_current_device_resource_ref().allocate(stream, 2 * sizeof(program));
93+
d_nodes1 = (node*)rmm::mr::get_current_device_resource_ref().allocate(
94+
stream, 7 * sizeof(node), alignof(node));
95+
d_nodes2 = (node*)rmm::mr::get_current_device_resource_ref().allocate(
96+
stream, 7 * sizeof(node), alignof(node));
97+
d_progs = (program_t)rmm::mr::get_current_device_resource_ref().allocate(
98+
stream, 2 * sizeof(program), alignof(program));
9799

98100
RAFT_CUDA_TRY(cudaMemcpyAsync(
99101
d_lYpred.data(), h_lYpred.data(), 500 * sizeof(float), cudaMemcpyHostToDevice, stream));
@@ -146,9 +148,12 @@ class GeneticProgramTest : public ::testing::Test {
146148

147149
void TearDown() override
148150
{
149-
rmm::mr::get_current_device_resource_ref().deallocate(stream, d_nodes1, 7 * sizeof(node));
150-
rmm::mr::get_current_device_resource_ref().deallocate(stream, d_nodes2, 7 * sizeof(node));
151-
rmm::mr::get_current_device_resource_ref().deallocate(stream, d_progs, 2 * sizeof(program));
151+
rmm::mr::get_current_device_resource_ref().deallocate(
152+
stream, d_nodes1, 7 * sizeof(node), alignof(node));
153+
rmm::mr::get_current_device_resource_ref().deallocate(
154+
stream, d_nodes2, 7 * sizeof(node), alignof(node));
155+
rmm::mr::get_current_device_resource_ref().deallocate(
156+
stream, d_progs, 2 * sizeof(program), alignof(program));
152157
}
153158

154159
raft::handle_t handle;

python/cuml/cuml/common/classification.py

Lines changed: 2 additions & 131 deletions
Original file line numberDiff line numberDiff line change
@@ -1,142 +1,13 @@
1-
# SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
1+
# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
22
# SPDX-License-Identifier: Apache-2.0
3-
import warnings
4-
53
import cudf
64
import cupy as cp
75
import numpy as np
8-
import pandas as pd
96

107
from cuml.internals.array import CumlArray, cuda_ptr
11-
from cuml.internals.input_utils import input_to_cuml_array, input_to_cupy_array
8+
from cuml.internals.input_utils import input_to_cupy_array
129
from cuml.internals.output_utils import cudf_to_pandas
1310

14-
is_integral = cp.ReductionKernel(
15-
"T x",
16-
"bool out",
17-
"ceilf(x) == x",
18-
"a && b",
19-
"out = a",
20-
"true",
21-
"is_integral",
22-
)
23-
24-
25-
def check_classification_targets(y):
26-
"""Check if `y` is composed of valid class labels"""
27-
if y.dtype.kind == "f" and not is_integral(y):
28-
raise ValueError(
29-
"Unknown label type: continuous. Maybe you are trying to fit a "
30-
"classifier, which expects discrete classes on a regression target "
31-
"with continuous values."
32-
)
33-
34-
35-
def preprocess_labels(
36-
y, dtype=None, order="C", n_samples=None, allow_multitarget=False
37-
):
38-
"""Preprocess the `y` input to a classifier.
39-
40-
Parameters
41-
----------
42-
y : array-like
43-
The labels for fitting, may be any type cuml supports as input.
44-
dtype : dtype, optional
45-
The output dtype to use for the encoded labels. If not provided,
46-
a data-dependent integral type will be used.
47-
order : {"C", "F"}, optional
48-
The array order to use for the encoded labels.
49-
n_samples : int, optional
50-
If provided, will raise an error if the number of samples in `y`
51-
doesn't match.
52-
allow_multitarget : bool, optional
53-
Whether to allow multi-target labels.
54-
55-
Returns
56-
-------
57-
y_encoded : cp.ndarray
58-
The labels, encoded as integers in [0, n_classes - 1].
59-
classes : np.ndarray or list[np.ndarray]
60-
The classes as a numpy array, or a list of numpy arrays if
61-
y is multi-target.
62-
"""
63-
# cudf may coerce the dtype, store the original so we can cast back later
64-
y_dtype = y.dtype if isinstance(y, np.ndarray) else None
65-
66-
# No cuda container supports all dtypes. Here we coerce to cupy when
67-
# possible, falling back to cudf Series/DataFrame otherwise.
68-
if isinstance(y, np.ndarray) and y.dtype.kind in "iufb":
69-
y = cp.asarray(y)
70-
elif isinstance(y, pd.DataFrame):
71-
y = cudf.DataFrame(y)
72-
elif isinstance(y, pd.Series):
73-
y = cudf.Series(y)
74-
elif not isinstance(y, (cp.ndarray, cudf.DataFrame, cudf.Series)):
75-
# Non-numeric dtype, always go through cudf
76-
y = input_to_cuml_array(y, convert_to_mem_type=False).array
77-
if y.dtype.kind in "iufb":
78-
y = y.to_output("cupy")
79-
else:
80-
y = (cudf.DataFrame if y.ndim == 2 else cudf.Series)(
81-
y, dtype=(np.dtype("O") if y.dtype.kind in "U" else None)
82-
)
83-
84-
# Validate dimensionality, ensuring 1D/2D y is as expected
85-
if y.ndim == 2 and y.shape[1] == 1:
86-
warnings.warn(
87-
"A column-vector y was passed when a 1d array was expected. Please "
88-
"change the shape of y to (n_samples,), for example using ravel()."
89-
)
90-
y = y.iloc[:, 0] if isinstance(y, cudf.DataFrame) else y.ravel()
91-
elif allow_multitarget and y.ndim not in (1, 2):
92-
raise ValueError(
93-
f"y should be a 1d or 2d array, got an array of shape {y.shape} instead."
94-
)
95-
elif not allow_multitarget and y.ndim != 1:
96-
raise ValueError(
97-
f"y should be a 1d array, got an array of shape {y.shape} instead."
98-
)
99-
100-
# Validate correct number of samples
101-
if n_samples is not None and y.shape[0] != n_samples:
102-
raise ValueError(
103-
f"Expected `y` with {n_samples} samples, got {y.shape[0]}"
104-
)
105-
106-
def _encode(y):
107-
"""Encode `y` to codes and classes"""
108-
check_classification_targets(y)
109-
if isinstance(y, cudf.Series):
110-
y = y.astype("category")
111-
codes = cp.asarray(y.cat.codes)
112-
classes = y.cat.categories.to_numpy()
113-
# cudf will sometimes translate non-numeric dtypes. Coerce back to
114-
# the input dtype if the input was originally a numpy array.
115-
if y_dtype is not None:
116-
classes = classes.astype(y_dtype, copy=False)
117-
else:
118-
classes, codes = cp.unique(y, return_inverse=True)
119-
classes = classes.get()
120-
return codes, classes
121-
122-
if y.ndim == 1:
123-
y_encoded, classes = _encode(y)
124-
if dtype is not None:
125-
y_encoded = y_encoded.astype(dtype, copy=False)
126-
else:
127-
getter = y.iloc if isinstance(y, cudf.DataFrame) else y
128-
encoded_cols, classes = zip(
129-
*(_encode(getter[:, i]) for i in range(y.shape[1]))
130-
)
131-
classes = list(classes)
132-
if dtype is None:
133-
dtype = cp.result_type(*(c.dtype for c in encoded_cols))
134-
y_encoded = cp.empty(shape=y.shape, dtype=dtype, order=order)
135-
for i, col in enumerate(encoded_cols):
136-
y_encoded[:, i] = col
137-
138-
return y_encoded, classes
139-
14011

14112
def decode_labels(y_encoded, classes, output_type="cupy"):
14213
"""Convert encoded labels back into their original classes.

python/cuml/cuml/ensemble/randomforestclassifier.py

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -7,14 +7,14 @@
77
import cuml.internals
88
import cuml.internals.nvtx as nvtx
99
from cuml.common.array_descriptor import CumlArrayDescriptor
10-
from cuml.common.classification import decode_labels, preprocess_labels
10+
from cuml.common.classification import decode_labels
1111
from cuml.common.doc_utils import generate_docstring, insert_into_docstring
1212
from cuml.ensemble.randomforest_common import BaseRandomForestModel
1313
from cuml.internals.array import CumlArray
1414
from cuml.internals.input_utils import input_to_cuml_array
1515
from cuml.internals.interop import UnsupportedOnGPU
1616
from cuml.internals.mixins import ClassifierMixin
17-
from cuml.internals.validation import check_features
17+
from cuml.internals.validation import check_features, check_y
1818
from cuml.metrics import accuracy_score
1919

2020

@@ -222,15 +222,14 @@ def fit(self, X, y, *, convert_dtype=True) -> "RandomForestClassifier":
222222
y to be of dtype int32. This will increase memory used for
223223
the method.
224224
"""
225+
y, classes = check_y(y, dtype=cp.int32, return_classes=True)
225226
X_m = input_to_cuml_array(
226227
X,
227228
convert_to_dtype=(np.float32 if convert_dtype else None),
228229
check_dtype=[np.float32, np.float64],
229230
order="F",
231+
check_rows=y.shape[0],
230232
).array
231-
y, classes = preprocess_labels(
232-
y, n_samples=X_m.shape[0], dtype=cp.int32
233-
)
234233
self.classes_ = classes
235234
self.n_classes_ = len(classes)
236235
y_m = CumlArray(data=y)

0 commit comments

Comments
 (0)