From 4a03d23bcf1526120e072516c3fc20c93b8932da Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Wed, 24 Apr 2024 06:42:41 -0700 Subject: [PATCH 01/15] added alpha parameter support to linear regression descriptor --- cpp/oneapi/dal/algo/linear_regression/common.cpp | 12 ++++++++++++ cpp/oneapi/dal/algo/linear_regression/common.hpp | 15 +++++++++++++++ 2 files changed, 27 insertions(+) diff --git a/cpp/oneapi/dal/algo/linear_regression/common.cpp b/cpp/oneapi/dal/algo/linear_regression/common.cpp index 70fd04f221e..92044817402 100644 --- a/cpp/oneapi/dal/algo/linear_regression/common.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/common.cpp @@ -42,6 +42,8 @@ class descriptor_impl : public base { explicit descriptor_impl() = default; bool compute_intercept = true; + double alpha = 0; + result_option_id result_options = get_default_result_options(); }; @@ -71,6 +73,16 @@ descriptor_base::descriptor_base(bool compute_intercept) impl_->compute_intercept = compute_intercept; } +template +double descriptor_base::get_alpha() const { + return impl_->alpha; +} + +template +void descriptor_base::set_alpha_impl(double value) { + impl_->alpha = value; +} + template bool descriptor_base::get_compute_intercept() const { return impl_->compute_intercept; diff --git a/cpp/oneapi/dal/algo/linear_regression/common.hpp b/cpp/oneapi/dal/algo/linear_regression/common.hpp index 633e919f1bb..9fc24344b31 100644 --- a/cpp/oneapi/dal/algo/linear_regression/common.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/common.hpp @@ -112,10 +112,12 @@ class descriptor_base : public base { descriptor_base(bool compute_intercept); bool get_compute_intercept() const; + double get_alpha() const; result_option_id get_result_options() const; protected: void set_compute_intercept_impl(bool compute_intercept); + void set_alpha_impl(double alpha); void set_result_options_impl(const result_option_id& value); private: @@ -165,6 +167,19 @@ class descriptor : public detail::descriptor_base { /// Creates a new instance of the class with default parameters explicit descriptor() : base_t(true) {} + explicit descriptor(bool compute_intercept, double alpha) : base_t(compute_intercept) { + set_alpha(alpha); + } + + explicit descriptor(double alpha) : base_t(true) { + set_alpha(alpha); + } + + auto& set_alpha(double value) { + base_t::set_alpha_impl(value); + return *this; + } + /// Defines should intercept be taken into consideration. bool get_compute_intercept() const { return base_t::get_compute_intercept(); From 3ef36b2614d5a207c2787715dd6e9e3f11550857 Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Thu, 2 May 2024 04:12:43 -0700 Subject: [PATCH 02/15] added support for rigde regression penalty in linear regression on batch gpu implementation --- .../backend/gpu/train_kernel_norm_eq_dpc.cpp | 28 +++++++++++++++++++ .../dal/algo/linear_regression/common.hpp | 2 +- 2 files changed, 29 insertions(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp index bf0cd04c00e..aaa536cbd68 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp @@ -37,6 +37,29 @@ using dal::backend::context_gpu; namespace be = dal::backend; namespace pr = be::primitives; +template +sycl::event add_ridge_penalty(sycl::queue& q, + pr::ndview& xtx, + double alpha, + const be::event_vector& deps) { + ONEDAL_ASSERT(xtx.has_mutable_data()); + ONEDAL_ASSERT(be::is_known_usm(q, xtx.get_mutable_data())); + ONEDAL_ASSERT(xtx.get_dimension(0) == xtx.get_dimension(1)); + + Float* xtx_ptr = xtx.get_mutable_data(); + std::int64_t feature_count = xtx.get_dimension(0); + + return q.submit([&](sycl::handler& cgh) { + const auto range = be::make_range_1d(feature_count); + cgh.depends_on(deps); + cgh.parallel_for(range, [=](sycl::id<1> idx) { + const std::int64_t diag_idx = + idx * (feature_count + 1); + xtx_ptr[diag_idx] += alpha; + }); + }); +} + template static train_result call_dal_kernel(const context_gpu& ctx, const detail::descriptor_base& desc, @@ -105,6 +128,11 @@ static train_result call_dal_kernel(const context_gpu& ctx, const be::event_vector solve_deps{ last_xty_event, last_xtx_event }; + double alpha = desc.get_alpha(); + if (alpha != 0) { + last_xtx_event = add_ridge_penalty(queue, xtx, desc.get_alpha(), { last_xtx_event }); + } + auto& comm = ctx.get_communicator(); if (comm.get_rank_count() > 1) { sycl::event::wait_and_throw(solve_deps); diff --git a/cpp/oneapi/dal/algo/linear_regression/common.hpp b/cpp/oneapi/dal/algo/linear_regression/common.hpp index 9fc24344b31..d0d27605206 100644 --- a/cpp/oneapi/dal/algo/linear_regression/common.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/common.hpp @@ -170,7 +170,7 @@ class descriptor : public detail::descriptor_base { explicit descriptor(bool compute_intercept, double alpha) : base_t(compute_intercept) { set_alpha(alpha); } - + explicit descriptor(double alpha) : base_t(true) { set_alpha(alpha); } From 64c8a82836031675e9f9a4504a5ddca7e2fffabc Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Wed, 8 May 2024 04:32:22 -0700 Subject: [PATCH 03/15] fixed gpu implementation, added switch to daal ridge on cpu --- .../algo/linear_regression/backend/cpu/BUILD | 1 + .../backend/cpu/train_kernel_norm_eq.cpp | 88 ++++++++++++------- .../backend/gpu/train_kernel_norm_eq_dpc.cpp | 22 ++--- 3 files changed, 69 insertions(+), 42 deletions(-) diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/BUILD b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/BUILD index 55adfee47a9..7bd3d6e679d 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/BUILD +++ b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/BUILD @@ -15,5 +15,6 @@ dal_module( "@onedal//cpp/daal:core", "@onedal//cpp/daal/src/algorithms/linear_model:kernel", "@onedal//cpp/daal/src/algorithms/linear_regression:kernel", + "@onedal//cpp/daal/src/algorithms/ridge_regression:kernel" ], ) diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp index dbea53a33f6..f27b22fed06 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp @@ -16,6 +16,7 @@ #include #include +#include #include "oneapi/dal/backend/interop/common.hpp" #include "oneapi/dal/backend/interop/error_converter.hpp" @@ -39,21 +40,26 @@ namespace be = dal::backend; namespace pr = be::primitives; namespace interop = dal::backend::interop; namespace daal_lr = daal::algorithms::linear_regression; +namespace daal_rr = daal::algorithms::ridge_regression; -using daal_hyperparameters_t = daal_lr::internal::Hyperparameter; +using daal_lr_hyperparameters_t = daal_lr::internal::Hyperparameter; -constexpr auto daal_method = daal_lr::training::normEqDense; +constexpr auto daal_lr_method = daal_lr::training::normEqDense; +constexpr auto daal_rr_method = daal_rr::training::normEqDense; template -using online_kernel_t = daal_lr::training::internal::OnlineKernel; +using online_lr_kernel_t = daal_lr::training::internal::OnlineKernel; + +template +using online_rr_kernel_t = daal_rr::training::internal::BatchKernel; template -static daal_hyperparameters_t convert_parameters(const detail::train_parameters& params) { +static daal_lr_hyperparameters_t convert_parameters(const detail::train_parameters& params) { using daal_lr::internal::HyperparameterId; const std::int64_t block = params.get_cpu_macro_block(); - daal_hyperparameters_t daal_hyperparameter; + daal_lr_hyperparameters_t daal_hyperparameter; auto status = daal_hyperparameter.set(HyperparameterId::denseUpdateStepBlockSize, block); interop::status_to_exception(status); @@ -97,33 +103,53 @@ static train_result call_daal_kernel(const context_cpu& ctx, auto x_daal_table = interop::convert_to_daal_table(data); auto y_daal_table = interop::convert_to_daal_table(resp); - const daal_hyperparameters_t& hp = convert_parameters(params); - - { - const auto status = interop::call_daal_kernel(ctx, - *x_daal_table, - *y_daal_table, - *xtx_daal_table, - *xty_daal_table, - intp, - &hp); - - interop::status_to_exception(status); + const daal_lr_hyperparameters_t& hp = convert_parameters(params); + + double alpha = desc.get_alpha(); + if (alpha != 0.0) { + auto ridge_matrix_array = array::full(1, static_cast(alpha)); + auto ridge_matrix = + interop::convert_to_daal_homogen_table(ridge_matrix_array, 1, 1); + + { + const auto status = interop::call_daal_kernel(ctx, + *x_daal_table, + *y_daal_table, + *xtx_daal_table, + *xty_daal_table, + *betas_daal_table, + intp, *ridge_matrix); + + interop::status_to_exception(status); + } } - - { - const auto status = dal::backend::dispatch_by_cpu(ctx, [&](auto cpu) { - constexpr auto cpu_type = interop::to_daal_cpu_type::value; - return online_kernel_t().finalizeCompute(*xtx_daal_table, - *xty_daal_table, - *xtx_daal_table, - *xty_daal_table, - *betas_daal_table, - intp, - &hp); - }); - - interop::status_to_exception(status); + else { + { + const auto status = interop::call_daal_kernel(ctx, + *x_daal_table, + *y_daal_table, + *xtx_daal_table, + *xty_daal_table, + intp, + &hp); + + interop::status_to_exception(status); + } + + { + const auto status = dal::backend::dispatch_by_cpu(ctx, [&](auto cpu) { + constexpr auto cpu_type = interop::to_daal_cpu_type::value; + return online_lr_kernel_t().finalizeCompute(*xtx_daal_table, + *xty_daal_table, + *xtx_daal_table, + *xty_daal_table, + *betas_daal_table, + intp, + &hp); + }); + + interop::status_to_exception(status); + } } auto betas_table = homogen_table::wrap(betas_arr, response_count, feature_count + 1); diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp index aaa536cbd68..1e4e0e7925b 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp @@ -40,6 +40,7 @@ namespace pr = be::primitives; template sycl::event add_ridge_penalty(sycl::queue& q, pr::ndview& xtx, + bool compute_intercept, double alpha, const be::event_vector& deps) { ONEDAL_ASSERT(xtx.has_mutable_data()); @@ -48,14 +49,13 @@ sycl::event add_ridge_penalty(sycl::queue& q, Float* xtx_ptr = xtx.get_mutable_data(); std::int64_t feature_count = xtx.get_dimension(0); + std::int64_t original_feature_count = feature_count - compute_intercept; return q.submit([&](sycl::handler& cgh) { - const auto range = be::make_range_1d(feature_count); + const auto range = be::make_range_1d(original_feature_count); cgh.depends_on(deps); cgh.parallel_for(range, [=](sycl::id<1> idx) { - const std::int64_t diag_idx = - idx * (feature_count + 1); - xtx_ptr[diag_idx] += alpha; + xtx_ptr[idx * (feature_count + 1)] += alpha; }); }); } @@ -85,8 +85,8 @@ static train_result call_dal_kernel(const context_gpu& ctx, const auto feature_count = data.get_column_count(); const auto response_count = resp.get_column_count(); ONEDAL_ASSERT(sample_count == resp.get_row_count()); - const bool beta = desc.get_compute_intercept(); - const std::int64_t ext_feature_count = feature_count + beta; + const bool compute_intercept = desc.get_compute_intercept(); + const std::int64_t ext_feature_count = feature_count + compute_intercept; const auto betas_size = check_mul_overflow(response_count, feature_count + 1); auto betas_arr = array::zeros(queue, betas_size, alloc); @@ -118,8 +118,8 @@ static train_result call_dal_kernel(const context_gpu& ctx, auto y_arr = y_accessor.pull(queue, { first, last }, alloc); auto y = pr::ndview::wrap(y_arr.get_data(), { length, response_count }); - last_xty_event = update_xty(queue, beta, x, y, xty, { last_xty_event }); - last_xtx_event = update_xtx(queue, beta, x, xtx, { last_xtx_event }); + last_xty_event = update_xty(queue, compute_intercept, x, y, xty, { last_xty_event }); + last_xtx_event = update_xtx(queue, compute_intercept, x, xtx, { last_xtx_event }); // We keep the latest slice of data up to date because of pimpl - // it virtually extend lifetime of pulled arrays @@ -129,8 +129,8 @@ static train_result call_dal_kernel(const context_gpu& ctx, const be::event_vector solve_deps{ last_xty_event, last_xtx_event }; double alpha = desc.get_alpha(); - if (alpha != 0) { - last_xtx_event = add_ridge_penalty(queue, xtx, desc.get_alpha(), { last_xtx_event }); + if (alpha != 0.0) { + last_xtx_event = add_ridge_penalty(queue, xtx, compute_intercept, alpha, { last_xtx_event }); } auto& comm = ctx.get_communicator(); @@ -150,7 +150,7 @@ static train_result call_dal_kernel(const context_gpu& ctx, auto nxtx = pr::ndarray::empty(queue, xtx_shape, alloc); auto nxty = pr::ndview::wrap_mutable(betas_arr, betas_shape); - auto solve_event = pr::solve_system(queue, beta, xtx, xty, nxtx, nxty, solve_deps); + auto solve_event = pr::solve_system(queue, compute_intercept, xtx, xty, nxtx, nxty, solve_deps); sycl::event::wait_and_throw({ solve_event }); auto betas = homogen_table::wrap(betas_arr, response_count, feature_count + 1); From cd3fb2ecea4629673909f599ab55a347a7675dbf Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Wed, 8 May 2024 06:06:00 -0700 Subject: [PATCH 04/15] added gpu/cpu implementations of ridge regression for online mode --- .../cpu/finalize_train_kernel_norm_eq.cpp | 70 +++++++++++++------ .../cpu/partial_train_kernel_norm_eq.cpp | 8 +-- .../backend/cpu/train_kernel_norm_eq.cpp | 54 +++++++------- .../gpu/finalize_train_kernel_norm_eq_dpc.cpp | 39 ++++++++++- .../gpu/partial_train_kernel_norm_eq_dpc.cpp | 16 +++-- .../backend/gpu/train_kernel_norm_eq_dpc.cpp | 8 ++- 6 files changed, 132 insertions(+), 63 deletions(-) diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/finalize_train_kernel_norm_eq.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/finalize_train_kernel_norm_eq.cpp index 5540641d8fd..7a40b46bf5d 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/finalize_train_kernel_norm_eq.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/finalize_train_kernel_norm_eq.cpp @@ -16,6 +16,7 @@ #include #include +#include #include "oneapi/dal/backend/interop/common.hpp" #include "oneapi/dal/backend/interop/error_converter.hpp" @@ -37,21 +38,26 @@ namespace be = dal::backend; namespace pr = be::primitives; namespace interop = dal::backend::interop; namespace daal_lr = daal::algorithms::linear_regression; +namespace daal_rr = daal::algorithms::ridge_regression; -using daal_hyperparameters_t = daal_lr::internal::Hyperparameter; +using daal_lr_hyperparameters_t = daal_lr::internal::Hyperparameter; -constexpr auto daal_method = daal_lr::training::normEqDense; +constexpr auto daal_lr_method = daal_lr::training::normEqDense; +constexpr auto daal_rr_method = daal_rr::training::normEqDense; template -using online_kernel_t = daal_lr::training::internal::OnlineKernel; +using online_lr_kernel_t = daal_lr::training::internal::OnlineKernel; + +template +using online_rr_kernel_t = daal_rr::training::internal::OnlineKernel; template -static daal_hyperparameters_t convert_parameters(const detail::train_parameters& params) { +static daal_lr_hyperparameters_t convert_parameters(const detail::train_parameters& params) { using daal_lr::internal::HyperparameterId; const std::int64_t block = params.get_cpu_macro_block(); - daal_hyperparameters_t daal_hyperparameter; + daal_lr_hyperparameters_t daal_hyperparameter; auto status = daal_hyperparameter.set(HyperparameterId::denseUpdateStepBlockSize, block); interop::status_to_exception(status); @@ -68,36 +74,58 @@ static train_result call_daal_kernel(const context_cpu& ctx, using model_t = model; using model_impl_t = detail::model_impl; - const bool beta = desc.get_compute_intercept(); + const bool compute_intercept = desc.get_compute_intercept(); const auto response_count = input.get_partial_xty().get_row_count(); const auto ext_feature_count = input.get_partial_xty().get_column_count(); - const auto feature_count = ext_feature_count - beta; + const auto feature_count = ext_feature_count - compute_intercept; const auto betas_size = check_mul_overflow(response_count, feature_count + 1); auto betas_arr = array::zeros(betas_size); - const daal_hyperparameters_t& hp = convert_parameters(params); + const daal_lr_hyperparameters_t& hp = convert_parameters(params); auto xtx_daal_table = interop::convert_to_daal_table(input.get_partial_xtx()); auto xty_daal_table = interop::convert_to_daal_table(input.get_partial_xty()); auto betas_daal_table = interop::convert_to_daal_homogen_table(betas_arr, response_count, feature_count + 1); - { - const auto status = dal::backend::dispatch_by_cpu(ctx, [&](auto cpu) { - constexpr auto cpu_type = interop::to_daal_cpu_type::value; - return online_kernel_t().finalizeCompute(*xtx_daal_table, - *xty_daal_table, - *xtx_daal_table, - *xty_daal_table, - *betas_daal_table, - beta, - &hp); - }); - - interop::status_to_exception(status); + double alpha = desc.get_alpha(); + if (alpha != 0.0) { + auto ridge_matrix_array = array::full(1, static_cast(alpha)); + auto ridge_matrix = interop::convert_to_daal_homogen_table(ridge_matrix_array, 1, 1); + + { + const auto status = dal::backend::dispatch_by_cpu(ctx, [&](auto cpu) { + constexpr auto cpu_type = interop::to_daal_cpu_type::value; + return online_rr_kernel_t().finalizeCompute(*xtx_daal_table, + *xty_daal_table, + *xtx_daal_table, + *xty_daal_table, + *betas_daal_table, + compute_intercept, + *ridge_matrix); + }); + + interop::status_to_exception(status); + } + } + else { + { + const auto status = dal::backend::dispatch_by_cpu(ctx, [&](auto cpu) { + constexpr auto cpu_type = interop::to_daal_cpu_type::value; + return online_lr_kernel_t().finalizeCompute(*xtx_daal_table, + *xty_daal_table, + *xtx_daal_table, + *xty_daal_table, + *betas_daal_table, + compute_intercept, + &hp); + }); + + interop::status_to_exception(status); + } } auto betas_table = homogen_table::wrap(betas_arr, response_count, feature_count + 1); diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/partial_train_kernel_norm_eq.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/partial_train_kernel_norm_eq.cpp index 7cac1aa47b7..d5d9f61003c 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/partial_train_kernel_norm_eq.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/partial_train_kernel_norm_eq.cpp @@ -62,14 +62,14 @@ static partial_train_result call_daal_kernel(const context_cpu& ctx, const partial_train_input& input) { using dal::detail::check_mul_overflow; - const bool beta = desc.get_compute_intercept(); + const bool compute_intercept = desc.get_compute_intercept(); const auto feature_count = input.get_data().get_column_count(); const auto response_count = input.get_responses().get_column_count(); const daal_hyperparameters_t& hp = convert_parameters(params); - const auto ext_feature_count = feature_count + beta; + const auto ext_feature_count = feature_count + compute_intercept; const bool has_xtx_data = input.get_prev().get_partial_xtx().has_data(); if (has_xtx_data) { @@ -85,7 +85,7 @@ static partial_train_result call_daal_kernel(const context_cpu& ctx, *y_daal_table, *daal_xtx, *daal_xty, - beta, + compute_intercept, &hp); interop::status_to_exception(status); @@ -117,7 +117,7 @@ static partial_train_result call_daal_kernel(const context_cpu& ctx, *y_daal_table, *xtx_daal_table, *xty_daal_table, - beta, + compute_intercept, &hp); interop::status_to_exception(status); diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp index f27b22fed06..eaa509d40cc 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp @@ -51,7 +51,7 @@ template using online_lr_kernel_t = daal_lr::training::internal::OnlineKernel; template -using online_rr_kernel_t = daal_rr::training::internal::BatchKernel; +using batch_rr_kernel_t = daal_rr::training::internal::BatchKernel; template static daal_lr_hyperparameters_t convert_parameters(const detail::train_parameters& params) { @@ -108,30 +108,32 @@ static train_result call_daal_kernel(const context_cpu& ctx, double alpha = desc.get_alpha(); if (alpha != 0.0) { auto ridge_matrix_array = array::full(1, static_cast(alpha)); - auto ridge_matrix = - interop::convert_to_daal_homogen_table(ridge_matrix_array, 1, 1); - - { - const auto status = interop::call_daal_kernel(ctx, - *x_daal_table, - *y_daal_table, - *xtx_daal_table, - *xty_daal_table, - *betas_daal_table, - intp, *ridge_matrix); + auto ridge_matrix = interop::convert_to_daal_homogen_table(ridge_matrix_array, 1, 1); + + { + const auto status = + interop::call_daal_kernel(ctx, + *x_daal_table, + *y_daal_table, + *xtx_daal_table, + *xty_daal_table, + *betas_daal_table, + intp, + *ridge_matrix); interop::status_to_exception(status); } } else { - { - const auto status = interop::call_daal_kernel(ctx, - *x_daal_table, - *y_daal_table, - *xtx_daal_table, - *xty_daal_table, - intp, - &hp); + { + const auto status = + interop::call_daal_kernel(ctx, + *x_daal_table, + *y_daal_table, + *xtx_daal_table, + *xty_daal_table, + intp, + &hp); interop::status_to_exception(status); } @@ -140,12 +142,12 @@ static train_result call_daal_kernel(const context_cpu& ctx, const auto status = dal::backend::dispatch_by_cpu(ctx, [&](auto cpu) { constexpr auto cpu_type = interop::to_daal_cpu_type::value; return online_lr_kernel_t().finalizeCompute(*xtx_daal_table, - *xty_daal_table, - *xtx_daal_table, - *xty_daal_table, - *betas_daal_table, - intp, - &hp); + *xty_daal_table, + *xtx_daal_table, + *xty_daal_table, + *betas_daal_table, + intp, + &hp); }); interop::status_to_exception(status); diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel_norm_eq_dpc.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel_norm_eq_dpc.cpp index 733bb46b0b3..3d1dc655761 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel_norm_eq_dpc.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel_norm_eq_dpc.cpp @@ -35,6 +35,27 @@ using dal::backend::context_gpu; namespace be = dal::backend; namespace pr = be::primitives; +template +sycl::event add_ridge_penalty(sycl::queue& q, + const pr::ndarray& xtx, + bool compute_intercept, + double alpha) { + ONEDAL_ASSERT(xtx.has_mutable_data()); + ONEDAL_ASSERT(be::is_known_usm(q, xtx.get_mutable_data())); + ONEDAL_ASSERT(xtx.get_dimension(0) == xtx.get_dimension(1)); + + Float* xtx_ptr = xtx.get_mutable_data(); + std::int64_t feature_count = xtx.get_dimension(0); + std::int64_t original_feature_count = feature_count - compute_intercept; + + return q.submit([&](sycl::handler& cgh) { + const auto range = be::make_range_1d(original_feature_count); + cgh.parallel_for(range, [=](sycl::id<1> idx) { + xtx_ptr[idx * (feature_count + 1)] += alpha; + }); + }); +} + template static train_result call_dal_kernel(const context_gpu& ctx, const detail::descriptor_base& desc, @@ -47,14 +68,14 @@ static train_result call_dal_kernel(const context_gpu& ctx, auto& queue = ctx.get_queue(); - const bool beta = desc.get_compute_intercept(); + const bool compute_intercept = desc.get_compute_intercept(); constexpr auto uplo = pr::mkl::uplo::upper; constexpr auto alloc = sycl::usm::alloc::device; const auto response_count = input.get_partial_xty().get_row_count(); const auto ext_feature_count = input.get_partial_xty().get_column_count(); - const auto feature_count = ext_feature_count - beta; + const auto feature_count = ext_feature_count - compute_intercept; const pr::ndshape<2> xtx_shape{ ext_feature_count, ext_feature_count }; @@ -69,9 +90,21 @@ static train_result call_dal_kernel(const context_gpu& ctx, const auto betas_size = check_mul_overflow(response_count, feature_count + 1); auto betas_arr = array::zeros(queue, betas_size, alloc); + double alpha = desc.get_alpha(); + sycl::event ridge_event; + if (alpha != 0.0) { + ridge_event = add_ridge_penalty(queue, xtx_nd, compute_intercept, alpha); + } + auto nxtx = pr::ndarray::empty(queue, xtx_shape, alloc); auto nxty = pr::ndview::wrap_mutable(betas_arr, betas_shape); - auto solve_event = pr::solve_system(queue, beta, xtx_nd, xty_nd, nxtx, nxty, {}); + auto solve_event = pr::solve_system(queue, + compute_intercept, + xtx_nd, + xty_nd, + nxtx, + nxty, + { ridge_event }); sycl::event::wait_and_throw({ solve_event }); auto betas = homogen_table::wrap(betas_arr, response_count, feature_count + 1); diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/partial_train_kernel_norm_eq_dpc.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/partial_train_kernel_norm_eq_dpc.cpp index dff0548afe4..a9aa7c373e4 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/partial_train_kernel_norm_eq_dpc.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/partial_train_kernel_norm_eq_dpc.cpp @@ -45,11 +45,11 @@ static partial_train_result call_dal_kernel(const context_gpu& ctx, constexpr auto alloc = sycl::usm::alloc::device; - const bool beta = desc.get_compute_intercept(); + const bool compute_intercept = desc.get_compute_intercept(); const auto feature_count = input.get_data().get_column_count(); const auto response_count = input.get_responses().get_column_count(); - const std::int64_t ext_feature_count = feature_count + beta; + const std::int64_t ext_feature_count = feature_count + compute_intercept; const pr::ndshape<2> xty_shape{ response_count, ext_feature_count }; const pr::ndshape<2> xtx_shape{ ext_feature_count, ext_feature_count }; @@ -74,8 +74,10 @@ static partial_train_result call_dal_kernel(const context_gpu& ctx, input_.get_partial_xty(), sycl::usm::alloc::device); auto copy_xty_event = copy(queue, xty, xty_nd, { fill_xty_event }); - auto last_xtx_event = update_xtx(queue, beta, data_nd, xtx, { copy_xtx_event }); - auto last_xty_event = update_xty(queue, beta, data_nd, res_nd, xty, { copy_xty_event }); + auto last_xtx_event = + update_xtx(queue, compute_intercept, data_nd, xtx, { copy_xtx_event }); + auto last_xty_event = + update_xty(queue, compute_intercept, data_nd, res_nd, xty, { copy_xty_event }); result.set_partial_xtx(homogen_table::wrap(xtx.flatten(queue, { last_xtx_event }), ext_feature_count, @@ -97,8 +99,10 @@ static partial_train_result call_dal_kernel(const context_gpu& ctx, auto [xtx, fill_xtx_event] = pr::ndarray::zeros(queue, xtx_shape, alloc); - auto last_xty_event = update_xty(queue, beta, data_nd, res_nd, xty, { fill_xty_event }); - auto last_xtx_event = update_xtx(queue, beta, data_nd, xtx, { fill_xtx_event }); + auto last_xty_event = + update_xty(queue, compute_intercept, data_nd, res_nd, xty, { fill_xty_event }); + auto last_xtx_event = + update_xtx(queue, compute_intercept, data_nd, xtx, { fill_xtx_event }); result.set_partial_xtx(homogen_table::wrap(xtx.flatten(queue, { last_xtx_event }), ext_feature_count, diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp index 1e4e0e7925b..39ace0d875c 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp @@ -40,7 +40,7 @@ namespace pr = be::primitives; template sycl::event add_ridge_penalty(sycl::queue& q, pr::ndview& xtx, - bool compute_intercept, + bool compute_intercept, double alpha, const be::event_vector& deps) { ONEDAL_ASSERT(xtx.has_mutable_data()); @@ -130,7 +130,8 @@ static train_result call_dal_kernel(const context_gpu& ctx, double alpha = desc.get_alpha(); if (alpha != 0.0) { - last_xtx_event = add_ridge_penalty(queue, xtx, compute_intercept, alpha, { last_xtx_event }); + last_xtx_event = + add_ridge_penalty(queue, xtx, compute_intercept, alpha, { last_xtx_event }); } auto& comm = ctx.get_communicator(); @@ -150,7 +151,8 @@ static train_result call_dal_kernel(const context_gpu& ctx, auto nxtx = pr::ndarray::empty(queue, xtx_shape, alloc); auto nxty = pr::ndview::wrap_mutable(betas_arr, betas_shape); - auto solve_event = pr::solve_system(queue, compute_intercept, xtx, xty, nxtx, nxty, solve_deps); + auto solve_event = + pr::solve_system(queue, compute_intercept, xtx, xty, nxtx, nxty, solve_deps); sycl::event::wait_and_throw({ solve_event }); auto betas = homogen_table::wrap(betas_arr, response_count, feature_count + 1); From f6706a60a902083a58c4c15a0231e060f2ecca9a Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Fri, 10 May 2024 01:27:27 -0700 Subject: [PATCH 05/15] added dll linking for BatchKernel on ridge regression from daal --- .../ridge_regression_train_dense_normeq_batch_fpt_cpu.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/daal/src/algorithms/ridge_regression/ridge_regression_train_dense_normeq_batch_fpt_cpu.cpp b/cpp/daal/src/algorithms/ridge_regression/ridge_regression_train_dense_normeq_batch_fpt_cpu.cpp index 2c71f4d64a0..e1ed3085861 100644 --- a/cpp/daal/src/algorithms/ridge_regression/ridge_regression_train_dense_normeq_batch_fpt_cpu.cpp +++ b/cpp/daal/src/algorithms/ridge_regression/ridge_regression_train_dense_normeq_batch_fpt_cpu.cpp @@ -41,7 +41,7 @@ template class BatchContainer; namespace internal { -template class BatchKernel; +template class DAAL_EXPORT BatchKernel; } // namespace internal } // namespace training From 0f39e211822fe6d74f9f414eefd5134b22f86f63 Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Fri, 10 May 2024 03:42:57 -0700 Subject: [PATCH 06/15] dll linking for online kernel --- .../ridge_regression_train_dense_normeq_online_fpt_cpu.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/daal/src/algorithms/ridge_regression/ridge_regression_train_dense_normeq_online_fpt_cpu.cpp b/cpp/daal/src/algorithms/ridge_regression/ridge_regression_train_dense_normeq_online_fpt_cpu.cpp index 867f3a23b56..c82553c834a 100644 --- a/cpp/daal/src/algorithms/ridge_regression/ridge_regression_train_dense_normeq_online_fpt_cpu.cpp +++ b/cpp/daal/src/algorithms/ridge_regression/ridge_regression_train_dense_normeq_online_fpt_cpu.cpp @@ -40,7 +40,7 @@ template class OnlineContainer; namespace internal { -template class OnlineKernel; +template class DAAL_EXPORT OnlineKernel; } // namespace internal } // namespace training From 3630121ce649ab498dad88aaad92b6a30aaf186a Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Mon, 13 May 2024 07:49:33 -0700 Subject: [PATCH 07/15] modified tests to support ridge regression in both batch and online modes --- .../dal/algo/linear_regression/test/batch.cpp | 10 +- .../algo/linear_regression/test/fixture.hpp | 135 +++++++++++++++++- .../algo/linear_regression/test/online.cpp | 10 +- .../dal/algo/linear_regression/test/spmd.cpp | 2 +- .../test/train_parameters.cpp | 2 +- 5 files changed, 151 insertions(+), 8 deletions(-) diff --git a/cpp/oneapi/dal/algo/linear_regression/test/batch.cpp b/cpp/oneapi/dal/algo/linear_regression/test/batch.cpp index 270b34b9ddc..00ec7babbb9 100644 --- a/cpp/oneapi/dal/algo/linear_regression/test/batch.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/test/batch.cpp @@ -47,7 +47,15 @@ TEMPLATE_LIST_TEST_M(lr_batch_test, "LR common flow", "[lr][batch]", lr_types) { this->generate(777); - this->run_and_check(); + this->run_and_check_linear(); +} + +TEMPLATE_LIST_TEST_M(lr_batch_test, "RR common flow", "[rr][batch]", lr_types) { + SKIP_IF(this->not_float64_friendly()); + + this->generate(777); + + this->run_and_check_ridge(); } } // namespace oneapi::dal::linear_regression::test diff --git a/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp b/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp index a8994a7c704..eae19361f2a 100644 --- a/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp @@ -123,6 +123,17 @@ class lr_test : public te::crtp_algo_fixture { return result; } + double generate_alpha(std::int64_t seed) const { + std::mt19937 gen(seed); + + double alpha_min = 0.1; + double alpha_max = 5; + + std::uniform_real_distribution dist(alpha_min, alpha_max); + + return dist(gen); + } + void check_table_dimensions(const table& x_train, const table& y_train, const table& x_test, @@ -144,13 +155,14 @@ class lr_test : public te::crtp_algo_fixture { this->bias_ = std::move(bias); this->beta_ = std::move(beta); + this->alpha_ = generate_alpha(seed); } - auto get_descriptor() const { + auto get_descriptor(double alpha = 0.0) const { result_option_id resopts = result_options::coefficients; if (this->intercept_) resopts = resopts | result_options::intercept; - return linear_regression::descriptor(intercept_) + return linear_regression::descriptor(intercept_, alpha) .set_result_options(resopts); } @@ -191,7 +203,63 @@ class lr_test : public te::crtp_algo_fixture { } } - void run_and_check(std::int64_t seed = 888, double tol = 1e-2) { + void check_coefficient_shrinkage(const table& lr_coeffs, + const table& rr_coeffs, + double tol = 1e-3) { + row_accessor lr_acc(lr_coeffs); + row_accessor rr_acc(rr_coeffs); + const auto lr_arr = lr_acc.pull({ 0, -1 }); + const auto rr_arr = rr_acc.pull({ 0, -1 }); + + double lr_norm = 0, rr_norm = 0; + for (std::int64_t i = 0; i < lr_arr.get_count(); ++i) { + lr_norm += lr_arr[i] * lr_arr[i]; + rr_norm += rr_arr[i] * rr_arr[i]; + } + lr_norm = std::sqrt(lr_norm); + rr_norm = std::sqrt(rr_norm); + + REQUIRE(rr_norm <= lr_norm + tol); + } + + void run_and_check_ridge(std::int64_t seed = 888, double tol = 1e-2) { + using namespace ::oneapi::dal::detail; + + std::mt19937 meta_gen(seed); + + const std::int64_t train_seed = meta_gen(); + const auto train_dataframe = GENERATE_DATAFRAME( + te::dataframe_builder{ this->s_count_, this->f_count_ }.fill_uniform(-5.5, + 3.5, + train_seed)); + auto x_train = train_dataframe.get_table(this->get_homogen_table_id()); + + const std::int64_t test_seed = meta_gen(); + const auto test_dataframe = GENERATE_DATAFRAME( + te::dataframe_builder{ this->t_count_, this->f_count_ }.fill_uniform(-3.5, + 5.5, + test_seed)); + auto x_test = test_dataframe.get_table(this->get_homogen_table_id()); + + auto y_train = compute_responses(this->beta_, this->bias_, x_train); + auto y_test = compute_responses(this->beta_, this->bias_, x_test); + + check_table_dimensions(x_train, y_train, x_test, y_test); + + const auto linear_desc = this->get_descriptor(); + const auto linear_train_res = this->train(linear_desc, x_train, y_train); + + const auto ridge_desc = this->get_descriptor(this->alpha_); + const auto ridge_train_res = this->train(ridge_desc, x_train, y_train); + + SECTION("Checking coefficient shrinkage") { + this->check_coefficient_shrinkage(linear_train_res.get_coefficients(), + ridge_train_res.get_coefficients(), + tol); + } + } + + void run_and_check_linear(std::int64_t seed = 888, double tol = 1e-2) { using namespace ::oneapi::dal::detail; std::mt19937 meta_gen(seed); @@ -234,6 +302,7 @@ class lr_test : public te::crtp_algo_fixture { check_if_close(infer_res.get_responses(), y_test, tol); } } + template std::vector split_table_by_rows(const dal::table& t, std::int64_t split_count) { ONEDAL_ASSERT(0l < split_count); @@ -259,7 +328,8 @@ class lr_test : public te::crtp_algo_fixture { return result; } - void run_and_check_online(std::int64_t nBlocks) { + + void run_and_check_linear_online(std::int64_t nBlocks) { using namespace ::oneapi::dal::detail; std::int64_t seed = 888; @@ -312,6 +382,62 @@ class lr_test : public te::crtp_algo_fixture { } } + void run_and_check_ridge_online(std::int64_t nBlocks) { + using namespace ::oneapi::dal::detail; + + std::int64_t seed = 888; + double tol = 1e-2; + + std::mt19937 meta_gen(seed); + const std::int64_t train_seed = meta_gen(); + const auto train_dataframe = GENERATE_DATAFRAME( + te::dataframe_builder{ this->s_count_, this->f_count_ }.fill_uniform(-5.5, + 3.5, + train_seed)); + auto x_train = train_dataframe.get_table(this->get_homogen_table_id()); + + const std::int64_t test_seed = meta_gen(); + const auto test_dataframe = GENERATE_DATAFRAME( + te::dataframe_builder{ this->t_count_, this->f_count_ }.fill_uniform(-3.5, + 5.5, + test_seed)); + auto x_test = test_dataframe.get_table(this->get_homogen_table_id()); + + auto y_train = compute_responses(this->beta_, this->bias_, x_train); + auto y_test = compute_responses(this->beta_, this->bias_, x_test); + + check_table_dimensions(x_train, y_train, x_test, y_test); + + auto input_table_x = split_table_by_rows(x_train, nBlocks); + auto input_table_y = split_table_by_rows(y_train, nBlocks); + + const auto linear_desc = this->get_descriptor(); + dal::linear_regression::partial_train_result<> linear_partial_result; + for (std::int64_t i = 0; i < nBlocks; i++) { + linear_partial_result = this->partial_train(linear_desc, + linear_partial_result, + input_table_x[i], + input_table_y[i]); + } + auto linear_train_res = this->finalize_train(linear_desc, linear_partial_result); + + const auto ridge_desc = this->get_descriptor(this->alpha_); + dal::linear_regression::partial_train_result<> ridge_partial_result; + for (std::int64_t i = 0; i < nBlocks; i++) { + ridge_partial_result = this->partial_train(ridge_desc, + ridge_partial_result, + input_table_x[i], + input_table_y[i]); + } + auto ridge_train_res = this->finalize_train(ridge_desc, ridge_partial_result); + + SECTION("Checking coefficient shrinkage") { + this->check_coefficient_shrinkage(linear_train_res.get_coefficients(), + ridge_train_res.get_coefficients(), + tol); + } + } + protected: bool intercept_ = true; std::int64_t t_count_; @@ -321,6 +447,7 @@ class lr_test : public te::crtp_algo_fixture { table bias_; table beta_; + double alpha_; }; using lr_types = COMBINE_TYPES((float, double), diff --git a/cpp/oneapi/dal/algo/linear_regression/test/online.cpp b/cpp/oneapi/dal/algo/linear_regression/test/online.cpp index 2724768491b..7a841a7e805 100644 --- a/cpp/oneapi/dal/algo/linear_regression/test/online.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/test/online.cpp @@ -47,7 +47,15 @@ TEMPLATE_LIST_TEST_M(lr_online_test, "LR common flow", "[lr][online]", lr_types) this->generate(777); const int64_t nBlocks = GENERATE(1, 3, 5, 8); - this->run_and_check_online(nBlocks); + this->run_and_check_linear_online(nBlocks); +} + +TEMPLATE_LIST_TEST_M(lr_online_test, "RR common flow", "[rr][online]", lr_types) { + SKIP_IF(this->not_float64_friendly()); + this->generate(777); + const int64_t nBlocks = GENERATE(1, 3, 5, 8); + + this->run_and_check_linear_online(nBlocks); } } // namespace oneapi::dal::linear_regression::test diff --git a/cpp/oneapi/dal/algo/linear_regression/test/spmd.cpp b/cpp/oneapi/dal/algo/linear_regression/test/spmd.cpp index d0cca4e943c..62223f03fdd 100644 --- a/cpp/oneapi/dal/algo/linear_regression/test/spmd.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/test/spmd.cpp @@ -25,7 +25,7 @@ TEMPLATE_LIST_TEST_M(lr_spmd_test, "LR common flow", "[lr][spmd]", lr_types) { this->generate(777); this->set_rank_count(GENERATE(2, 3)); - this->run_and_check(); + this->run_and_check_linear(); } } // namespace oneapi::dal::linear_regression::test diff --git a/cpp/oneapi/dal/algo/linear_regression/test/train_parameters.cpp b/cpp/oneapi/dal/algo/linear_regression/test/train_parameters.cpp index 48f9ead5d3a..835b8ecc1b4 100644 --- a/cpp/oneapi/dal/algo/linear_regression/test/train_parameters.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/test/train_parameters.cpp @@ -89,7 +89,7 @@ TEMPLATE_LIST_TEST_M(lr_train_params_test, "LR train params", "[lr][train][param this->generate(999); this->generate_parameters(); - this->run_and_check(); + this->run_and_check_linear(); } } // namespace oneapi::dal::linear_regression::test From 872c7dd3adbfa2882652dbac672121e8866b9174 Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Tue, 14 May 2024 03:29:59 -0700 Subject: [PATCH 08/15] refactored code per suggestions and adjusted tests to reduce floating point inaccuracies --- .../backend/gpu/train_kernel_norm_eq_dpc.cpp | 3 ++- cpp/oneapi/dal/algo/linear_regression/common.cpp | 16 ++++++++-------- cpp/oneapi/dal/algo/linear_regression/common.hpp | 14 +++++++++----- .../dal/algo/linear_regression/test/fixture.hpp | 12 +++++------- 4 files changed, 24 insertions(+), 21 deletions(-) diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp index 39ace0d875c..09ab536afd0 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp @@ -54,8 +54,9 @@ sycl::event add_ridge_penalty(sycl::queue& q, return q.submit([&](sycl::handler& cgh) { const auto range = be::make_range_1d(original_feature_count); cgh.depends_on(deps); + std::int64_t step = feature_count + 1; cgh.parallel_for(range, [=](sycl::id<1> idx) { - xtx_ptr[idx * (feature_count + 1)] += alpha; + xtx_ptr[idx * step] += alpha; }); }); } diff --git a/cpp/oneapi/dal/algo/linear_regression/common.cpp b/cpp/oneapi/dal/algo/linear_regression/common.cpp index 92044817402..949898f3524 100644 --- a/cpp/oneapi/dal/algo/linear_regression/common.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/common.cpp @@ -74,23 +74,23 @@ descriptor_base::descriptor_base(bool compute_intercept) } template -double descriptor_base::get_alpha() const { - return impl_->alpha; +bool descriptor_base::get_compute_intercept() const { + return impl_->compute_intercept; } template -void descriptor_base::set_alpha_impl(double value) { - impl_->alpha = value; +void descriptor_base::set_compute_intercept_impl(bool compute_intercept) { + impl_->compute_intercept = compute_intercept; } template -bool descriptor_base::get_compute_intercept() const { - return impl_->compute_intercept; +double descriptor_base::get_alpha() const { + return impl_->alpha; } template -void descriptor_base::set_compute_intercept_impl(bool compute_intercept) { - impl_->compute_intercept = compute_intercept; +void descriptor_base::set_alpha_impl(double value) { + impl_->alpha = value; } template class ONEDAL_EXPORT descriptor_base; diff --git a/cpp/oneapi/dal/algo/linear_regression/common.hpp b/cpp/oneapi/dal/algo/linear_regression/common.hpp index d0d27605206..9b7affcd138 100644 --- a/cpp/oneapi/dal/algo/linear_regression/common.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/common.hpp @@ -175,11 +175,6 @@ class descriptor : public detail::descriptor_base { set_alpha(alpha); } - auto& set_alpha(double value) { - base_t::set_alpha_impl(value); - return *this; - } - /// Defines should intercept be taken into consideration. bool get_compute_intercept() const { return base_t::get_compute_intercept(); @@ -190,6 +185,15 @@ class descriptor : public detail::descriptor_base { return *this; } + double get_alpha() const { + return base_t::get_alpha(); + } + + auto& set_alpha(double value) { + base_t::set_alpha_impl(value); + return *this; + } + /// Choose which results should be computed and returned. result_option_id get_result_options() const { return base_t::get_result_options(); diff --git a/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp b/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp index eae19361f2a..14f37a72f53 100644 --- a/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp @@ -211,15 +211,13 @@ class lr_test : public te::crtp_algo_fixture { const auto lr_arr = lr_acc.pull({ 0, -1 }); const auto rr_arr = rr_acc.pull({ 0, -1 }); - double lr_norm = 0, rr_norm = 0; + double lr_norm_squared = 0, rr_norm_squared = 0; for (std::int64_t i = 0; i < lr_arr.get_count(); ++i) { - lr_norm += lr_arr[i] * lr_arr[i]; - rr_norm += rr_arr[i] * rr_arr[i]; + lr_norm_squared += lr_arr[i] * lr_arr[i]; + rr_norm_squared += rr_arr[i] * rr_arr[i]; } - lr_norm = std::sqrt(lr_norm); - rr_norm = std::sqrt(rr_norm); - REQUIRE(rr_norm <= lr_norm + tol); + REQUIRE(rr_norm_squared <= lr_norm_squared + tol); } void run_and_check_ridge(std::int64_t seed = 888, double tol = 1e-2) { @@ -440,6 +438,7 @@ class lr_test : public te::crtp_algo_fixture { protected: bool intercept_ = true; + double alpha_; std::int64_t t_count_; std::int64_t s_count_; std::int64_t f_count_; @@ -447,7 +446,6 @@ class lr_test : public te::crtp_algo_fixture { table bias_; table beta_; - double alpha_; }; using lr_types = COMBINE_TYPES((float, double), From a49208c006ceb6eb4d57859e704decc601d00fbf Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Tue, 14 May 2024 05:28:48 -0700 Subject: [PATCH 09/15] generalized adding ridge penalty for both batch and online modes --- .../gpu/finalize_train_kernel_norm_eq_dpc.cpp | 22 +------ .../linear_regression/backend/gpu/misc.hpp | 66 +++++++++++++++++++ .../backend/gpu/train_kernel_norm_eq_dpc.cpp | 25 +------ .../algo/linear_regression/test/fixture.hpp | 2 +- 4 files changed, 69 insertions(+), 46 deletions(-) create mode 100644 cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel_norm_eq_dpc.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel_norm_eq_dpc.cpp index 3d1dc655761..d3431663249 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel_norm_eq_dpc.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel_norm_eq_dpc.cpp @@ -27,6 +27,7 @@ #include "oneapi/dal/algo/linear_regression/backend/model_impl.hpp" #include "oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel.hpp" #include "oneapi/dal/algo/linear_regression/backend/gpu/update_kernel.hpp" +#include "oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp" namespace oneapi::dal::linear_regression::backend { @@ -35,27 +36,6 @@ using dal::backend::context_gpu; namespace be = dal::backend; namespace pr = be::primitives; -template -sycl::event add_ridge_penalty(sycl::queue& q, - const pr::ndarray& xtx, - bool compute_intercept, - double alpha) { - ONEDAL_ASSERT(xtx.has_mutable_data()); - ONEDAL_ASSERT(be::is_known_usm(q, xtx.get_mutable_data())); - ONEDAL_ASSERT(xtx.get_dimension(0) == xtx.get_dimension(1)); - - Float* xtx_ptr = xtx.get_mutable_data(); - std::int64_t feature_count = xtx.get_dimension(0); - std::int64_t original_feature_count = feature_count - compute_intercept; - - return q.submit([&](sycl::handler& cgh) { - const auto range = be::make_range_1d(original_feature_count); - cgh.parallel_for(range, [=](sycl::id<1> idx) { - xtx_ptr[idx * (feature_count + 1)] += alpha; - }); - }); -} - template static train_result call_dal_kernel(const context_gpu& ctx, const detail::descriptor_base& desc, diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp new file mode 100644 index 00000000000..e7015ed38d1 --- /dev/null +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp @@ -0,0 +1,66 @@ +/******************************************************************************* +* Copyright 2024 Intel Corporation +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +#include "oneapi/dal/detail/profiler.hpp" +#include "oneapi/dal/backend/primitives/ndarray.hpp" + +namespace oneapi::dal::linear_regression::backend { + +#ifdef ONEDAL_DATA_PARALLEL + +using alloc = sycl::usm::alloc; +namespace bk = dal::backend; +namespace pr = dal::backend::primitives; + +/// Adds ridge penalty to the diagonal elements of the xtx matrix + +/// +/// @tparam Float Floating-point type used to perform computations +/// +/// @param[in] q The SYCL queue +/// @param[in] xtx The input matrix to which the ridge penalty is added +/// @param[in] compute_intercept Flag indicating whether the intercept term is used in the matrix +/// @param[in] alpha The regularization parameter +/// @param[in] deps Events indicating the availability of the `xtx` for reading or writing +/// +/// @return A SYCL event indicating the availability of the matrix for reading and writing +template +sycl::event add_ridge_penalty(sycl::queue& q, + const pr::ndarray& xtx, + bool compute_intercept, + double alpha, + const bk::event_vector& deps = {}) { + ONEDAL_ASSERT(xtx.has_mutable_data()); + ONEDAL_ASSERT(be::is_known_usm(q, xtx.get_mutable_data())); + ONEDAL_ASSERT(xtx.get_dimension(0) == xtx.get_dimension(1)); + + Float* xtx_ptr = xtx.get_mutable_data(); + std::int64_t feature_count = xtx.get_dimension(0); + std::int64_t original_feature_count = feature_count - compute_intercept; + + return q.submit([&](sycl::handler& cgh) { + const auto range = be::make_range_1d(original_feature_count); + cgh.depends_on(deps); + std::int64_t step = feature_count + 1; + cgh.parallel_for(range, [=](sycl::id<1> idx) { + xtx_ptr[idx * step] += alpha; + }); + }); +} + +} // namespace oneapi::dal::linear_regression::backend + +#endif // ONEDAL_DATA_PARALLEL diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp index 09ab536afd0..25b08aa7710 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/train_kernel_norm_eq_dpc.cpp @@ -29,6 +29,7 @@ #include "oneapi/dal/algo/linear_regression/backend/model_impl.hpp" #include "oneapi/dal/algo/linear_regression/backend/gpu/train_kernel.hpp" #include "oneapi/dal/algo/linear_regression/backend/gpu/update_kernel.hpp" +#include "oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp" namespace oneapi::dal::linear_regression::backend { @@ -37,30 +38,6 @@ using dal::backend::context_gpu; namespace be = dal::backend; namespace pr = be::primitives; -template -sycl::event add_ridge_penalty(sycl::queue& q, - pr::ndview& xtx, - bool compute_intercept, - double alpha, - const be::event_vector& deps) { - ONEDAL_ASSERT(xtx.has_mutable_data()); - ONEDAL_ASSERT(be::is_known_usm(q, xtx.get_mutable_data())); - ONEDAL_ASSERT(xtx.get_dimension(0) == xtx.get_dimension(1)); - - Float* xtx_ptr = xtx.get_mutable_data(); - std::int64_t feature_count = xtx.get_dimension(0); - std::int64_t original_feature_count = feature_count - compute_intercept; - - return q.submit([&](sycl::handler& cgh) { - const auto range = be::make_range_1d(original_feature_count); - cgh.depends_on(deps); - std::int64_t step = feature_count + 1; - cgh.parallel_for(range, [=](sycl::id<1> idx) { - xtx_ptr[idx * step] += alpha; - }); - }); -} - template static train_result call_dal_kernel(const context_gpu& ctx, const detail::descriptor_base& desc, diff --git a/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp b/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp index 14f37a72f53..1bb18346c9b 100644 --- a/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp @@ -126,7 +126,7 @@ class lr_test : public te::crtp_algo_fixture { double generate_alpha(std::int64_t seed) const { std::mt19937 gen(seed); - double alpha_min = 0.1; + double alpha_min = 1; double alpha_max = 5; std::uniform_real_distribution dist(alpha_min, alpha_max); From 2b1f1e9351862c11fec9c271e5f05023b3df2519 Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Wed, 15 May 2024 01:35:07 -0700 Subject: [PATCH 10/15] fixed alpha float type for gpus not supporting fp64 --- cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp | 2 +- cpp/oneapi/dal/algo/linear_regression/test/online.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp b/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp index 1bb18346c9b..d263a061a22 100644 --- a/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp @@ -438,7 +438,7 @@ class lr_test : public te::crtp_algo_fixture { protected: bool intercept_ = true; - double alpha_; + float_t alpha_; std::int64_t t_count_; std::int64_t s_count_; std::int64_t f_count_; diff --git a/cpp/oneapi/dal/algo/linear_regression/test/online.cpp b/cpp/oneapi/dal/algo/linear_regression/test/online.cpp index 7a841a7e805..c16e1c06f26 100644 --- a/cpp/oneapi/dal/algo/linear_regression/test/online.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/test/online.cpp @@ -55,7 +55,7 @@ TEMPLATE_LIST_TEST_M(lr_online_test, "RR common flow", "[rr][online]", lr_types) this->generate(777); const int64_t nBlocks = GENERATE(1, 3, 5, 8); - this->run_and_check_linear_online(nBlocks); + this->run_and_check_ridge_online(nBlocks); } } // namespace oneapi::dal::linear_regression::test From 10274174ed27d316100dac50d3d2da9bd67dba41 Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Wed, 15 May 2024 02:42:42 -0700 Subject: [PATCH 11/15] adjusted alpha floating type from double to supported float for gpus not supporting fp64 --- cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp index e7015ed38d1..c81f8847014 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp @@ -41,7 +41,7 @@ template sycl::event add_ridge_penalty(sycl::queue& q, const pr::ndarray& xtx, bool compute_intercept, - double alpha, + Float alpha, const bk::event_vector& deps = {}) { ONEDAL_ASSERT(xtx.has_mutable_data()); ONEDAL_ASSERT(be::is_known_usm(q, xtx.get_mutable_data())); From ed6feaa5e4874d3123114841a00a98a16242a347 Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Thu, 16 May 2024 06:17:13 -0700 Subject: [PATCH 12/15] switched linreg batch cpu switch to use daal batch instead of online mode, minor optimizations and code refactoring --- .../cpu/finalize_train_kernel_norm_eq.cpp | 4 +- .../backend/cpu/train_kernel_norm_eq.cpp | 36 +++------ .../linear_regression/backend/gpu/misc.hpp | 8 +- .../dal/algo/linear_regression/common.hpp | 1 + .../algo/linear_regression/test/fixture.hpp | 81 ++++--------------- 5 files changed, 32 insertions(+), 98 deletions(-) diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/finalize_train_kernel_norm_eq.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/finalize_train_kernel_norm_eq.cpp index 7a40b46bf5d..88b1c58ccc4 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/finalize_train_kernel_norm_eq.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/finalize_train_kernel_norm_eq.cpp @@ -84,8 +84,6 @@ static train_result call_daal_kernel(const context_cpu& ctx, const auto betas_size = check_mul_overflow(response_count, feature_count + 1); auto betas_arr = array::zeros(betas_size); - const daal_lr_hyperparameters_t& hp = convert_parameters(params); - auto xtx_daal_table = interop::convert_to_daal_table(input.get_partial_xtx()); auto xty_daal_table = interop::convert_to_daal_table(input.get_partial_xty()); auto betas_daal_table = @@ -112,6 +110,8 @@ static train_result call_daal_kernel(const context_cpu& ctx, } } else { + const daal_lr_hyperparameters_t& hp = convert_parameters(params); + { const auto status = dal::backend::dispatch_by_cpu(ctx, [&](auto cpu) { constexpr auto cpu_type = interop::to_daal_cpu_type::value; diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp index eaa509d40cc..0e6e1f8cd10 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/cpu/train_kernel_norm_eq.cpp @@ -48,7 +48,7 @@ constexpr auto daal_lr_method = daal_lr::training::normEqDense; constexpr auto daal_rr_method = daal_rr::training::normEqDense; template -using online_lr_kernel_t = daal_lr::training::internal::OnlineKernel; +using batch_lr_kernel_t = daal_lr::training::internal::BatchKernel; template using batch_rr_kernel_t = daal_rr::training::internal::BatchKernel; @@ -103,8 +103,6 @@ static train_result call_daal_kernel(const context_cpu& ctx, auto x_daal_table = interop::convert_to_daal_table(data); auto y_daal_table = interop::convert_to_daal_table(resp); - const daal_lr_hyperparameters_t& hp = convert_parameters(params); - double alpha = desc.get_alpha(); if (alpha != 0.0) { auto ridge_matrix_array = array::full(1, static_cast(alpha)); @@ -125,30 +123,18 @@ static train_result call_daal_kernel(const context_cpu& ctx, } } else { - { - const auto status = - interop::call_daal_kernel(ctx, - *x_daal_table, - *y_daal_table, - *xtx_daal_table, - *xty_daal_table, - intp, - &hp); - - interop::status_to_exception(status); - } + const daal_lr_hyperparameters_t& hp = convert_parameters(params); { - const auto status = dal::backend::dispatch_by_cpu(ctx, [&](auto cpu) { - constexpr auto cpu_type = interop::to_daal_cpu_type::value; - return online_lr_kernel_t().finalizeCompute(*xtx_daal_table, - *xty_daal_table, - *xtx_daal_table, - *xty_daal_table, - *betas_daal_table, - intp, - &hp); - }); + const auto status = + interop::call_daal_kernel(ctx, + *x_daal_table, + *y_daal_table, + *xtx_daal_table, + *xty_daal_table, + *betas_daal_table, + intp, + &hp); interop::status_to_exception(status); } diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp index c81f8847014..ab65e196e6e 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp @@ -30,11 +30,11 @@ namespace pr = dal::backend::primitives; /// /// @tparam Float Floating-point type used to perform computations /// -/// @param[in] q The SYCL queue -/// @param[in] xtx The input matrix to which the ridge penalty is added +/// @param[in] q The SYCL queue +/// @param[in] xtx The input matrix to which the ridge penalty is added /// @param[in] compute_intercept Flag indicating whether the intercept term is used in the matrix -/// @param[in] alpha The regularization parameter -/// @param[in] deps Events indicating the availability of the `xtx` for reading or writing +/// @param[in] alpha The regularization parameter +/// @param[in] deps Events indicating the availability of the `xtx` for reading or writing /// /// @return A SYCL event indicating the availability of the matrix for reading and writing template diff --git a/cpp/oneapi/dal/algo/linear_regression/common.hpp b/cpp/oneapi/dal/algo/linear_regression/common.hpp index 9b7affcd138..57d597a984d 100644 --- a/cpp/oneapi/dal/algo/linear_regression/common.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/common.hpp @@ -185,6 +185,7 @@ class descriptor : public detail::descriptor_base { return *this; } + /// Defines regularization term alpha used in Ridge Regression double get_alpha() const { return base_t::get_alpha(); } diff --git a/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp b/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp index d263a061a22..aedf0165454 100644 --- a/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp @@ -220,7 +220,8 @@ class lr_test : public te::crtp_algo_fixture { REQUIRE(rr_norm_squared <= lr_norm_squared + tol); } - void run_and_check_ridge(std::int64_t seed = 888, double tol = 1e-2) { + std::tuple prepare_inputs(std::int64_t seed = 888, + double tol = 1e-2) { using namespace ::oneapi::dal::detail; std::mt19937 meta_gen(seed); @@ -243,6 +244,12 @@ class lr_test : public te::crtp_algo_fixture { auto y_test = compute_responses(this->beta_, this->bias_, x_test); check_table_dimensions(x_train, y_train, x_test, y_test); + return { x_train, y_train, x_test, y_test }; + } + + void run_and_check_ridge(std::int64_t seed = 888, double tol = 1e-2) { + table x_train, y_train, x_test, y_test; + std::tie(x_train, y_train, x_test, y_test) = prepare_inputs(seed, tol); const auto linear_desc = this->get_descriptor(); const auto linear_train_res = this->train(linear_desc, x_train, y_train); @@ -258,28 +265,8 @@ class lr_test : public te::crtp_algo_fixture { } void run_and_check_linear(std::int64_t seed = 888, double tol = 1e-2) { - using namespace ::oneapi::dal::detail; - - std::mt19937 meta_gen(seed); - - const std::int64_t train_seed = meta_gen(); - const auto train_dataframe = GENERATE_DATAFRAME( - te::dataframe_builder{ this->s_count_, this->f_count_ }.fill_uniform(-5.5, - 3.5, - train_seed)); - auto x_train = train_dataframe.get_table(this->get_homogen_table_id()); - - const std::int64_t test_seed = meta_gen(); - const auto test_dataframe = GENERATE_DATAFRAME( - te::dataframe_builder{ this->t_count_, this->f_count_ }.fill_uniform(-3.5, - 5.5, - test_seed)); - auto x_test = test_dataframe.get_table(this->get_homogen_table_id()); - - auto y_train = compute_responses(this->beta_, this->bias_, x_train); - auto y_test = compute_responses(this->beta_, this->bias_, x_test); - - check_table_dimensions(x_train, y_train, x_test, y_test); + table x_train, y_train, x_test, y_test; + std::tie(x_train, y_train, x_test, y_test) = prepare_inputs(seed, tol); const auto desc = this->get_descriptor(); const auto train_res = this->train(desc, x_train, y_train); @@ -328,30 +315,10 @@ class lr_test : public te::crtp_algo_fixture { } void run_and_check_linear_online(std::int64_t nBlocks) { - using namespace ::oneapi::dal::detail; - std::int64_t seed = 888; double tol = 1e-2; - - std::mt19937 meta_gen(seed); - const std::int64_t train_seed = meta_gen(); - const auto train_dataframe = GENERATE_DATAFRAME( - te::dataframe_builder{ this->s_count_, this->f_count_ }.fill_uniform(-5.5, - 3.5, - train_seed)); - auto x_train = train_dataframe.get_table(this->get_homogen_table_id()); - - const std::int64_t test_seed = meta_gen(); - const auto test_dataframe = GENERATE_DATAFRAME( - te::dataframe_builder{ this->t_count_, this->f_count_ }.fill_uniform(-3.5, - 5.5, - test_seed)); - auto x_test = test_dataframe.get_table(this->get_homogen_table_id()); - - auto y_train = compute_responses(this->beta_, this->bias_, x_train); - auto y_test = compute_responses(this->beta_, this->bias_, x_test); - - check_table_dimensions(x_train, y_train, x_test, y_test); + table x_train, y_train, x_test, y_test; + std::tie(x_train, y_train, x_test, y_test) = prepare_inputs(seed, tol); const auto desc = this->get_descriptor(); dal::linear_regression::partial_train_result<> partial_result; @@ -381,30 +348,10 @@ class lr_test : public te::crtp_algo_fixture { } void run_and_check_ridge_online(std::int64_t nBlocks) { - using namespace ::oneapi::dal::detail; - std::int64_t seed = 888; double tol = 1e-2; - - std::mt19937 meta_gen(seed); - const std::int64_t train_seed = meta_gen(); - const auto train_dataframe = GENERATE_DATAFRAME( - te::dataframe_builder{ this->s_count_, this->f_count_ }.fill_uniform(-5.5, - 3.5, - train_seed)); - auto x_train = train_dataframe.get_table(this->get_homogen_table_id()); - - const std::int64_t test_seed = meta_gen(); - const auto test_dataframe = GENERATE_DATAFRAME( - te::dataframe_builder{ this->t_count_, this->f_count_ }.fill_uniform(-3.5, - 5.5, - test_seed)); - auto x_test = test_dataframe.get_table(this->get_homogen_table_id()); - - auto y_train = compute_responses(this->beta_, this->bias_, x_train); - auto y_test = compute_responses(this->beta_, this->bias_, x_test); - - check_table_dimensions(x_train, y_train, x_test, y_test); + table x_train, y_train, x_test, y_test; + std::tie(x_train, y_train, x_test, y_test) = prepare_inputs(seed, tol); auto input_table_x = split_table_by_rows(x_train, nBlocks); auto input_table_y = split_table_by_rows(y_train, nBlocks); From d99b79eece2ffa5f83bef1e53df90ff47eab465d Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Fri, 17 May 2024 03:16:53 -0700 Subject: [PATCH 13/15] added dependency description of xtx on compute_intercept --- cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp index ab65e196e6e..ad4f17156be 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp @@ -32,7 +32,7 @@ namespace pr = dal::backend::primitives; /// /// @param[in] q The SYCL queue /// @param[in] xtx The input matrix to which the ridge penalty is added -/// @param[in] compute_intercept Flag indicating whether the intercept term is used in the matrix +/// @param[in] compute_intercept Flag indicating whether the intercept term is used in the matrix, extending it with extra dimension if true /// @param[in] alpha The regularization parameter /// @param[in] deps Events indicating the availability of the `xtx` for reading or writing /// From a1f030db48fbea321369f12dd92a8d910131834a Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Fri, 17 May 2024 05:42:39 -0700 Subject: [PATCH 14/15] fixed dll for daal lin reg batch kernel --- .../linear_regression_train_dense_normeq_batch_fpt_cpu.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/daal/src/algorithms/linear_regression/linear_regression_train_dense_normeq_batch_fpt_cpu.cpp b/cpp/daal/src/algorithms/linear_regression/linear_regression_train_dense_normeq_batch_fpt_cpu.cpp index ef9ab58c256..bb0a22d089e 100644 --- a/cpp/daal/src/algorithms/linear_regression/linear_regression_train_dense_normeq_batch_fpt_cpu.cpp +++ b/cpp/daal/src/algorithms/linear_regression/linear_regression_train_dense_normeq_batch_fpt_cpu.cpp @@ -39,7 +39,7 @@ template class BatchContainer; } namespace internal { -template class BatchKernel; +template class DAAL_EXPORT BatchKernel; } } // namespace training } // namespace linear_regression From 43f06e7a9de42f7d4259d03f5ff3a41ef336de6f Mon Sep 17 00:00:00 2001 From: Khalil Asadzade Date: Tue, 21 May 2024 02:28:58 -0700 Subject: [PATCH 15/15] updated copyright message for misc.hpp --- cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp index ad4f17156be..5ad5ba647ec 100644 --- a/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp +++ b/cpp/oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2024 Intel Corporation +* Copyright contributors to the oneDAL project * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License.