11/*
2- * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION.
2+ * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved .
33 * SPDX-License-Identifier: Apache-2.0
44 */
55
1515#include < raft/linalg/eig.cuh>
1616#include < raft/linalg/eltwise.cuh>
1717#include < raft/linalg/gemm.cuh>
18+ #include < raft/linalg/map.cuh>
1819#include < raft/linalg/pca_types.hpp>
20+ #include < raft/linalg/reduce.cuh>
1921#include < raft/linalg/rsvd.cuh>
2022#include < raft/linalg/transpose.cuh>
2123#include < raft/matrix/copy.cuh>
@@ -104,13 +106,19 @@ void cal_comp_exp_vars_svd(raft::resources const& handle,
104106 handle, explained_vars.data_handle (), explained_var_ratio.data_handle (), n_components, stream);
105107}
106108
107- template <typename math_t , typename idx_t >
109+ template <typename math_t , typename idx_t , typename LayoutPolicy >
108110void cal_eig (raft::resources const & handle,
109111 const paramsTSVD& prms,
110- raft::device_matrix_view<math_t , idx_t , raft::col_major > in,
111- raft::device_matrix_view<math_t , idx_t , raft::col_major > components,
112+ raft::device_matrix_view<math_t , idx_t , LayoutPolicy > in,
113+ raft::device_matrix_view<math_t , idx_t , LayoutPolicy > components,
112114 raft::device_vector_view<math_t , idx_t > explained_var)
113115{
116+ static_assert (
117+ std::is_same_v<LayoutPolicy, raft::row_major> || std::is_same_v<LayoutPolicy, raft::col_major>,
118+ " cal_eig: layout must be raft::row_major or raft::col_major" );
119+
120+ constexpr bool is_row_major = std::is_same_v<LayoutPolicy, raft::row_major>;
121+
114122 auto stream = resource::get_cuda_stream (handle);
115123 auto cusolver_handle = raft::resource::get_cusolver_dn_handle (handle);
116124
@@ -141,61 +149,68 @@ void cal_eig(raft::resources const& handle,
141149 raft::matrix::col_reverse (handle_stream_zero,
142150 raft::make_device_matrix_view<math_t , idx_t , raft::col_major>(
143151 components.data_handle (), n_cols, n_cols));
144- raft::linalg::transpose (components.data_handle (), n_cols, stream);
152+ if constexpr (!is_row_major) {
153+ raft::linalg::transpose (components.data_handle (), n_cols, stream);
154+ }
145155
146156 raft::matrix::row_reverse (handle_stream_zero,
147157 raft::make_device_matrix_view<math_t , idx_t , raft::row_major>(
148158 explained_var.data_handle (), n_cols, idx_t (1 )));
149159}
150160
151161/* *
152- * @brief sign flip for PCA and tSVD. Stabilizes the sign of column major eigenvectors.
162+ * @brief sign flip for PCA and tSVD. Stabilizes the sign of the eigenvectors.
163+ * @tparam math_t element type
164+ * @tparam idx_t index type
165+ * @tparam LayoutPolicy layout of the input and components matrices
153166 * @param handle: raft::resources
154- * @param input: input data [n_samples x n_features] (col-major)
155- * @param components: components matrix [n_components x n_features] (col-major)
167+ * @param input: input data [n_samples x n_features]
168+ * @param components: components matrix [n_components x n_features]
156169 * @param center whether to mean-center input before computing signs
157170 * @param flip_signs_based_on_U whether to determine signs by U (true) or V.T (false)
158171 */
159- template <typename math_t , typename idx_t >
172+ template <typename math_t , typename idx_t , typename LayoutPolicy >
160173void sign_flip_components (raft::resources const & handle,
161- raft::device_matrix_view<math_t , idx_t , raft::col_major > input,
162- raft::device_matrix_view<math_t , idx_t , raft::col_major > components,
174+ raft::device_matrix_view<math_t , idx_t , LayoutPolicy > input,
175+ raft::device_matrix_view<math_t , idx_t , LayoutPolicy > components,
163176 bool center,
164177 bool flip_signs_based_on_U = false )
165178{
179+ static_assert (
180+ std::is_same_v<LayoutPolicy, raft::row_major> || std::is_same_v<LayoutPolicy, raft::col_major>,
181+ " sign_flip_components: layout must be raft::row_major or raft::col_major" );
182+ constexpr bool is_row_major = std::is_same_v<LayoutPolicy, raft::row_major>;
183+
166184 auto stream = resource::get_cuda_stream (handle);
167185 auto n_samples = input.extent (0 );
168186 auto n_features = input.extent (1 );
169187 auto n_components = components.extent (0 );
170188
171189 rmm::device_uvector<math_t > max_vals (static_cast <std::size_t >(n_components), stream);
172- auto components_view = raft::make_device_matrix_view<math_t , idx_t , raft::col_major >(
190+ auto components_view = raft::make_device_matrix_view<math_t , idx_t , LayoutPolicy >(
173191 components.data_handle (), n_components, n_features);
174192 auto max_vals_view = raft::make_device_vector_view<math_t , idx_t >(max_vals.data (), n_components);
175193
176194 if (flip_signs_based_on_U) {
177195 if (center) {
178196 rmm::device_uvector<math_t > col_means (static_cast <std::size_t >(n_features), stream);
179- raft::stats::mean<false >(
197+ raft::stats::mean<is_row_major >(
180198 col_means.data (), input.data_handle (), n_features, n_samples, stream);
181- raft::stats::meanCenter<false , true >(
199+ raft::stats::meanCenter<is_row_major , true >(
182200 input.data_handle (), input.data_handle (), col_means.data (), n_features, n_samples, stream);
183201 }
184202 rmm::device_uvector<math_t > US (static_cast <std::size_t >(n_samples * n_components), stream);
185- raft::linalg::gemm<math_t , math_t , math_t , math_t >(handle,
186- input.data_handle (),
187- n_samples,
188- n_features,
189- components.data_handle (),
190- US .data (),
191- n_samples,
192- n_components,
193- CUBLAS_OP_N ,
194- CUBLAS_OP_T ,
195- math_t (1 ),
196- math_t (0 ),
197- stream);
198- raft::linalg::reduce<false , false >(
203+ raft::linalg::gemm (handle,
204+ input,
205+ raft::make_device_matrix_view<
206+ math_t ,
207+ idx_t ,
208+ std::conditional_t <is_row_major, raft::col_major, raft::row_major>>(
209+ components.data_handle (), n_features, n_components),
210+ raft::make_device_matrix_view<math_t , idx_t , LayoutPolicy>(
211+ US .data (), n_samples, n_components));
212+
213+ raft::linalg::reduce<is_row_major, false >(
199214 max_vals.data (),
200215 US .data (),
201216 n_components,
@@ -211,7 +226,7 @@ void sign_flip_components(raft::resources const& handle,
211226 },
212227 raft::identity_op ());
213228 } else {
214- raft::linalg::reduce<false , true >(
229+ raft::linalg::reduce<is_row_major , true >(
215230 max_vals.data (),
216231 components.data_handle (),
217232 n_features,
@@ -232,8 +247,15 @@ void sign_flip_components(raft::resources const& handle,
232247 handle,
233248 components_view,
234249 [components_view, max_vals_view, n_components, n_features] __device__ (auto idx) {
235- auto row = idx % n_components;
236- auto column = idx / n_components;
250+ idx_t row;
251+ idx_t column;
252+ if constexpr (is_row_major) {
253+ row = idx / n_features;
254+ column = idx % n_features;
255+ } else {
256+ row = idx % n_components;
257+ column = idx / n_components;
258+ }
237259 return (max_vals_view (row) < math_t (0 )) ? (-components_view (row, column))
238260 : components_view (row, column);
239261 });
@@ -376,18 +398,21 @@ void tsvd_fit(raft::resources const& handle,
376398 * @brief performs transform operation for the tsvd. Transforms the data to eigenspace.
377399 * @param[in] handle raft::resources
378400 * @param[in] prms: data structure that includes all the parameters from input size to algorithm.
379- * @param[in] input: the data to transform. Size n_rows x n_cols (col-major) .
380- * @param[in] components: principal components. Size n_components x n_cols (col-major) .
381- * @param[out] trans_input: transformed output. Size n_rows x n_components (col-major) .
401+ * @param[in] input: the data to transform. Size n_rows x n_cols.
402+ * @param[in] components: principal components. Size n_components x n_cols.
403+ * @param[out] trans_input: transformed output. Size n_rows x n_components.
382404 */
383- template <typename math_t , typename idx_t >
405+ template <typename math_t , typename idx_t , typename LayoutPolicy = raft::col_major >
384406void tsvd_transform (raft::resources const & handle,
385407 const paramsTSVD& prms,
386- raft::device_matrix_view<math_t , idx_t , raft::col_major > input,
387- raft::device_matrix_view<math_t , idx_t , raft::col_major > components,
388- raft::device_matrix_view<math_t , idx_t , raft::col_major > trans_input)
408+ raft::device_matrix_view<math_t , idx_t , LayoutPolicy > input,
409+ raft::device_matrix_view<math_t , idx_t , LayoutPolicy > components,
410+ raft::device_matrix_view<math_t , idx_t , LayoutPolicy > trans_input)
389411{
390- auto stream = resource::get_cuda_stream (handle);
412+ static_assert (
413+ std::is_same_v<LayoutPolicy, raft::row_major> || std::is_same_v<LayoutPolicy, raft::col_major>,
414+ " tsvd_transform: layout must be raft::row_major or raft::col_major" );
415+ constexpr bool is_row_major = std::is_same_v<LayoutPolicy, raft::row_major>;
391416
392417 auto n_rows = input.extent (0 );
393418 auto n_cols = input.extent (1 );
@@ -397,39 +422,34 @@ void tsvd_transform(raft::resources const& handle,
397422 ASSERT (n_rows > 0 , " Parameter n_rows: number of rows cannot be less than one" );
398423 ASSERT (n_components > 0 , " Parameter n_components: number of components cannot be less than one" );
399424
400- math_t alpha = math_t (1 );
401- math_t beta = math_t (0 );
402425 raft::linalg::gemm (handle,
403- input.data_handle (),
404- n_rows,
405- n_cols,
406- components.data_handle (),
407- trans_input.data_handle (),
408- n_rows,
409- n_components,
410- CUBLAS_OP_N ,
411- CUBLAS_OP_T ,
412- alpha,
413- beta,
414- stream);
426+ input,
427+ raft::make_device_matrix_view<
428+ math_t ,
429+ idx_t ,
430+ std::conditional_t <is_row_major, raft::col_major, raft::row_major>>(
431+ components.data_handle (), n_cols, n_components),
432+ trans_input);
415433}
416434
417435/* *
418436 * @brief performs inverse transform operation for the tsvd.
419437 * @param[in] handle raft::resources
420438 * @param[in] prms: data structure that includes all the parameters from input size to algorithm.
421- * @param[in] trans_input: the transformed data. Size n_rows x n_components (col-major) .
422- * @param[in] components: principal components. Size n_components x n_cols (col-major) .
423- * @param[out] output: reconstructed output. Size n_rows x n_cols (col-major) .
439+ * @param[in] trans_input: the transformed data. Size n_rows x n_components.
440+ * @param[in] components: principal components. Size n_components x n_cols.
441+ * @param[out] output: reconstructed output. Size n_rows x n_cols.
424442 */
425- template <typename math_t , typename idx_t >
443+ template <typename math_t , typename idx_t , typename LayoutPolicy = raft::col_major >
426444void tsvd_inverse_transform (raft::resources const & handle,
427445 const paramsTSVD& prms,
428- raft::device_matrix_view<math_t , idx_t , raft::col_major > trans_input,
429- raft::device_matrix_view<math_t , idx_t , raft::col_major > components,
430- raft::device_matrix_view<math_t , idx_t , raft::col_major > output)
446+ raft::device_matrix_view<math_t , idx_t , LayoutPolicy > trans_input,
447+ raft::device_matrix_view<math_t , idx_t , LayoutPolicy > components,
448+ raft::device_matrix_view<math_t , idx_t , LayoutPolicy > output)
431449{
432- auto stream = resource::get_cuda_stream (handle);
450+ static_assert (
451+ std::is_same_v<LayoutPolicy, raft::row_major> || std::is_same_v<LayoutPolicy, raft::col_major>,
452+ " tsvd_inverse_transform: layout must be raft::row_major or raft::col_major" );
433453
434454 auto n_rows = output.extent (0 );
435455 auto n_cols = output.extent (1 );
@@ -439,22 +459,7 @@ void tsvd_inverse_transform(raft::resources const& handle,
439459 ASSERT (n_rows > 0 , " Parameter n_rows: number of rows cannot be less than one" );
440460 ASSERT (n_components > 0 , " Parameter n_components: number of components cannot be less than one" );
441461
442- math_t alpha = math_t (1 );
443- math_t beta = math_t (0 );
444-
445- raft::linalg::gemm (handle,
446- trans_input.data_handle (),
447- n_rows,
448- n_components,
449- components.data_handle (),
450- output.data_handle (),
451- n_rows,
452- n_cols,
453- CUBLAS_OP_N ,
454- CUBLAS_OP_N ,
455- alpha,
456- beta,
457- stream);
462+ raft::linalg::gemm (handle, trans_input, components, output);
458463}
459464
460465/* *
0 commit comments