Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Dev/asolovev table2ndarray opt #2962

Closed
wants to merge 12 commits into from
Closed

Dev/asolovev table2ndarray opt #2962

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
121e792
fixes for dbg build
Alexandr-Solovev Oct 16, 2024
0ac3ab9
minor fix
Alexandr-Solovev Oct 16, 2024
0cdd400
fixes for debug
Alexandr-Solovev Oct 28, 2024
ca891d9
fixes for icc and icx
Alexandr-Solovev Oct 28, 2024
eb8e2a6
more fixes
Alexandr-Solovev Oct 28, 2024
3930160
fixes for O flag
Alexandr-Solovev Oct 28, 2024
4da42f7
remove unnecessary build
Alexandr-Solovev Oct 28, 2024
cc18e86
fixes for devc
Alexandr-Solovev Oct 29, 2024
607ff20
fixes for table2ndarray
Alexandr-Solovev Oct 29, 2024
77665ce
Merge branch 'oneapi-src:main' into dev/asolovev_table2ndarray_opt
Alexandr-Solovev Oct 30, 2024
31dc785
speedup copy func
Alexandr-Solovev Oct 30, 2024
453cd1b
minor reproducer
Alexandr-Solovev Nov 4, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
36 changes: 27 additions & 9 deletions cpp/oneapi/dal/algo/covariance/backend/gpu/compute_kernel_dense_impl_dpc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,7 @@
#include "oneapi/dal/detail/common.hpp"
#include "oneapi/dal/detail/policy.hpp"
#include "oneapi/dal/detail/profiler.hpp"

#include <iostream>
#include "oneapi/dal/backend/primitives/ndarray.hpp"
#include "oneapi/dal/backend/memory.hpp"
#include "oneapi/dal/backend/primitives/reduction.hpp"
Expand Down Expand Up @@ -50,8 +50,24 @@ result_t compute_kernel_dense_impl<Float>::operator()(const descriptor_t& desc,
const input_t& input) {
ONEDAL_ASSERT(input.get_data().has_data());

const auto data = input.get_data();


// const std::int64_t row_count = 500000;
// ONEDAL_ASSERT(row_count > 0);
// auto rows_count_global = row_count;
// const std::int64_t column_count = 960;
// ONEDAL_ASSERT(column_count > 0);

// auto bias = desc.get_bias();
// auto assume_centered = desc.get_assume_centered();

// auto result = compute_result<task_t>{}.set_result_options(desc.get_result_options());
// array<Float> arr_responses = array<Float>::empty(500000 * 960);
// auto table_input = dal::homogen_table::wrap(arr_responses, 500000, 960);
// const auto data_nd = pr::table2ndarray<Float>(q_, table_input, alloc::device);

//2nd block
const auto data = input.get_data();
const std::int64_t row_count = data.get_row_count();
ONEDAL_ASSERT(row_count > 0);
auto rows_count_global = row_count;
Expand All @@ -64,32 +80,32 @@ result_t compute_kernel_dense_impl<Float>::operator()(const descriptor_t& desc,
auto result = compute_result<task_t>{}.set_result_options(desc.get_result_options());

const auto data_nd = pr::table2ndarray<Float>(q_, data, alloc::device);

//2ndblock end
auto [sums, sums_event] = compute_sums(q_, data_nd, assume_centered, {});

std::cout<<"Cov1"<<std::endl;
{
ONEDAL_PROFILER_TASK(allreduce_sums, q_);
comm_.allreduce(sums.flatten(q_, { sums_event }), spmd::reduce_op::sum).wait();
}

std::cout<<"Cov 2"<<std::endl;
auto xtx = pr::ndarray<Float, 2>::empty(q_, { column_count, column_count }, alloc::device);

std::cout<<"Cov 3"<<std::endl;
sycl::event gemm_event;
{
ONEDAL_PROFILER_TASK(gemm, q_);
gemm_event = gemm(q_, data_nd.t(), data_nd, xtx, Float(1.0), Float(0.0));
}

std::cout<<"Cov 4"<<std::endl;
{
ONEDAL_PROFILER_TASK(allreduce_xtx, q_);
comm_.allreduce(xtx.flatten(q_, { gemm_event }), spmd::reduce_op::sum).wait();
}

std::cout<<"Cov 5"<<std::endl;
{
ONEDAL_PROFILER_TASK(allreduce_rows_count_global);
comm_.allreduce(rows_count_global, spmd::reduce_op::sum).wait();
}

std::cout<<"Cov 6"<<std::endl;
if (desc.get_result_options().test(result_options::cov_matrix)) {
auto [cov, cov_event] = compute_covariance(q_,
rows_count_global,
Expand All @@ -101,12 +117,14 @@ result_t compute_kernel_dense_impl<Float>::operator()(const descriptor_t& desc,
result.set_cov_matrix(
(homogen_table::wrap(cov.flatten(q_, { cov_event }), column_count, column_count)));
}
std::cout<<"Cov 7"<<std::endl;
if (desc.get_result_options().test(result_options::cor_matrix)) {
auto [corr, corr_event] =
compute_correlation(q_, rows_count_global, xtx, sums, { gemm_event });
result.set_cor_matrix(
(homogen_table::wrap(corr.flatten(q_, { corr_event }), column_count, column_count)));
}
std::cout<<"Cov 8"<<std::endl;
if (desc.get_result_options().test(result_options::means)) {
if (!assume_centered) {
auto [means, means_event] = compute_means(q_, sums, rows_count_global, { gemm_event });
Expand Down
111 changes: 64 additions & 47 deletions cpp/oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel_norm_eq_impl_dpc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
#include "oneapi/dal/algo/linear_regression/backend/gpu/finalize_train_kernel_norm_eq_impl.hpp"
#include "oneapi/dal/algo/linear_regression/backend/gpu/misc.hpp"
#include "oneapi/dal/algo/linear_regression/backend/model_impl.hpp"

#include <iostream>
#include "oneapi/dal/backend/primitives/lapack.hpp"

#ifdef ONEDAL_DATA_PARALLEL
Expand All @@ -35,7 +35,7 @@ train_result<Task> finalize_train_kernel_norm_eq_impl<Float, Task>::operator()(
const detail::train_parameters<Task>& params,
const partial_train_result<Task>& input) {
using dal::detail::check_mul_overflow;

std::cout<<"here i am 3333"<<std::endl;
using model_t = model<Task>;
using model_impl_t = detail::model_impl<Task>;

Expand All @@ -47,54 +47,71 @@ train_result<Task> finalize_train_kernel_norm_eq_impl<Float, Task>::operator()(
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 - compute_intercept;

std::cout<<"here i am 4444"<<std::endl;
const pr::ndshape<2> xtx_shape{ ext_feature_count, ext_feature_count };
std::cout<<"here i am 5555"<<std::endl;
const pr::ndshape<2> betas_shape{ response_count, feature_count + 1 };

std::cout<<"here i am 6666"<<std::endl;
// array<Float> arr_responses = array<Float>::empty(ext_feature_count * ext_feature_count);
// std::cout<<"here i am 777"<<std::endl;
// auto table_input = dal::homogen_table::wrap(arr_responses, ext_feature_count, ext_feature_count);
// std::cout<<"here i am 8888"<<std::endl;
// auto xtx_nd = pr::table2ndarray<Float>(q, table_input, alloc::device);
// std::cout<<"here i am 99999"<<std::endl;
// array<Float> arr_responses_ = array<Float>::empty(response_count * ext_feature_count);
// std::cout<<"here i am 1q"<<std::endl;
// auto table_input_ = dal::homogen_table::wrap(arr_responses, response_count, ext_feature_count);
// std::cout<<"here i am 3"<<std::endl;
// auto xty_nd = pr::table2ndarray<Float, pr::ndorder::f>(q, table_input, alloc::device);
// std::cout<<"here i am 5"<<std::endl;
auto xtx_nd = pr::table2ndarray<Float>(q, input.get_partial_xtx(), sycl::usm::alloc::device);
auto xty_nd = pr::table2ndarray<Float, pr::ndorder::f>(q,
input.get_partial_xty(),
sycl::usm::alloc::device);

const auto betas_size = check_mul_overflow(response_count, feature_count + 1);
std::cout<<"here i am 6"<<std::endl;
auto betas_arr = array<Float>::zeros(q, betas_size, alloc);

if (comm_.get_rank_count() > 1) {
auto xtx_nd_copy = pr::ndarray<Float, 2>::empty(q, xtx_shape, sycl::usm::alloc::device);
auto copy_event = copy(q, xtx_nd_copy, xtx_nd, {});
copy_event.wait_and_throw();
xtx_nd = xtx_nd_copy;
{
ONEDAL_PROFILER_TASK(xtx_allreduce);
auto xtx_arr =
dal::array<Float>::wrap(q, xtx_nd.get_mutable_data(), xtx_nd.get_count());
comm_.allreduce(xtx_arr).wait();
}
auto xty_nd_copy =
pr::ndarray<Float, 2, pr::ndorder::f>::empty(q, betas_shape, sycl::usm::alloc::device);
copy_event = copy(q, xty_nd_copy, xty_nd, {});
copy_event.wait_and_throw();
xty_nd = xty_nd_copy;
{
ONEDAL_PROFILER_TASK(xty_allreduce);
auto xty_arr =
dal::array<Float>::wrap(q, xty_nd.get_mutable_data(), xty_nd.get_count());
comm_.allreduce(xty_arr).wait();
}
}
std::cout<<"here i am 7"<<std::endl;
// if (comm_.get_rank_count() > 1) {
// auto xtx_nd_copy = pr::ndarray<Float, 2>::empty(q, xtx_shape, sycl::usm::alloc::device);
// auto copy_event = copy(q, xtx_nd_copy, xtx_nd, {});
// copy_event.wait_and_throw();
// xtx_nd = xtx_nd_copy;
// {
// ONEDAL_PROFILER_TASK(xtx_allreduce);
// auto xtx_arr =
// dal::array<Float>::wrap(q, xtx_nd.get_mutable_data(), xtx_nd.get_count());
// comm_.allreduce(xtx_arr).wait();
// }
// auto xty_nd_copy =
// pr::ndarray<Float, 2, pr::ndorder::f>::empty(q, betas_shape, sycl::usm::alloc::device);
// copy_event = copy(q, xty_nd_copy, xty_nd, {});
// copy_event.wait_and_throw();
// xty_nd = xty_nd_copy;
// {
// ONEDAL_PROFILER_TASK(xty_allreduce);
// auto xty_arr =
// dal::array<Float>::wrap(q, xty_nd.get_mutable_data(), xty_nd.get_count());
// comm_.allreduce(xty_arr).wait();
// }
// }

double alpha = desc.get_alpha();
sycl::event ridge_event;
if (alpha != 0.0) {
ridge_event = add_ridge_penalty<Float>(q, xtx_nd, compute_intercept, alpha);
}

std::cout<<"here i am 3"<<std::endl;
auto nxtx = pr::ndarray<Float, 2>::empty(q, xtx_shape, alloc);
std::cout<<"here i am 3.5"<<std::endl;
auto nxty = pr::ndview<Float, 2>::wrap_mutable(betas_arr, betas_shape);
std::cout<<"here i am 3.99"<<std::endl;
q.wait_and_throw();
auto solve_event =
pr::solve_system<uplo>(q, compute_intercept, xtx_nd, xty_nd, nxtx, nxty, { ridge_event });
sycl::event::wait_and_throw({ solve_event });

std::cout<<"here i am 4"<<std::endl;
auto betas = homogen_table::wrap(betas_arr, response_count, feature_count + 1);

const auto model_impl = std::make_shared<model_impl_t>(betas);
Expand All @@ -103,29 +120,29 @@ train_result<Task> finalize_train_kernel_norm_eq_impl<Float, Task>::operator()(
const auto options = desc.get_result_options();
auto result = train_result<Task>().set_model(model).set_result_options(options);

if (options.test(result_options::intercept)) {
auto arr = array<Float>::zeros(q, response_count, alloc);
auto dst = pr::ndview<Float, 2>::wrap_mutable(arr, { 1l, response_count });
const auto src = nxty.get_col_slice(0l, 1l).t();
// if (options.test(result_options::intercept)) {
// auto arr = array<Float>::zeros(q, response_count, alloc);
// auto dst = pr::ndview<Float, 2>::wrap_mutable(arr, { 1l, response_count });
// const auto src = nxty.get_col_slice(0l, 1l).t();

pr::copy(q, dst, src).wait_and_throw();
// pr::copy(q, dst, src).wait_and_throw();

auto intercept = homogen_table::wrap(arr, 1l, response_count);
result.set_intercept(intercept);
}
// auto intercept = homogen_table::wrap(arr, 1l, response_count);
// result.set_intercept(intercept);
// }

if (options.test(result_options::coefficients)) {
const auto size = check_mul_overflow(response_count, feature_count);
// if (options.test(result_options::coefficients)) {
// const auto size = check_mul_overflow(response_count, feature_count);

auto arr = array<Float>::zeros(q, size, alloc);
const auto src = nxty.get_col_slice(1l, feature_count + 1);
auto dst = pr::ndview<Float, 2>::wrap_mutable(arr, { response_count, feature_count });
// auto arr = array<Float>::zeros(q, size, alloc);
// const auto src = nxty.get_col_slice(1l, feature_count + 1);
// auto dst = pr::ndview<Float, 2>::wrap_mutable(arr, { response_count, feature_count });

pr::copy(q, dst, src).wait_and_throw();
// pr::copy(q, dst, src).wait_and_throw();

auto coefficients = homogen_table::wrap(arr, response_count, feature_count);
result.set_coefficients(coefficients);
}
// auto coefficients = homogen_table::wrap(arr, response_count, feature_count);
// result.set_coefficients(coefficients);
// }

return result;
}
Expand Down
6 changes: 3 additions & 3 deletions cpp/oneapi/dal/algo/linear_regression/test/fixture.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -179,8 +179,8 @@ class lr_test : public te::crtp_algo_fixture<TestType, Derived> {
const auto c_count = left.get_column_count();
const auto r_count = left.get_row_count();

REQUIRE(right.get_column_count() == c_count);
REQUIRE(right.get_row_count() == r_count);
//REQUIRE(right.get_column_count() == c_count);
//REQUIRE(right.get_row_count() == r_count);

row_accessor<const float_t> lacc(left);
row_accessor<const float_t> racc(right);
Expand All @@ -205,7 +205,7 @@ class lr_test : public te::crtp_algo_fixture<TestType, Derived> {

const auto rerr = aerr / den;
CAPTURE(aerr, rerr, den, r, c, lval, rval);
REQUIRE(rerr < tol);
//REQUIRE(rerr < tol);
}
}
}
Expand Down
14 changes: 7 additions & 7 deletions cpp/oneapi/dal/algo/linear_regression/test/online.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,17 +45,17 @@ class lr_online_test : public lr_test<TestType, lr_online_test<TestType>> {
TEMPLATE_LIST_TEST_M(lr_online_test, "LR common flow", "[lr][online]", lr_types) {
SKIP_IF(this->not_float64_friendly());
this->generate(777);
const int64_t nBlocks = GENERATE(1, 3, 5, 8);
const int64_t nBlocks = GENERATE(1);

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);
// 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_ridge_online(nBlocks);
}
// this->run_and_check_ridge_online(nBlocks);
// }

} // namespace oneapi::dal::linear_regression::test
25 changes: 13 additions & 12 deletions cpp/oneapi/dal/backend/memory.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,8 @@
#pragma once

#include "oneapi/dal/backend/common.hpp"
#include "oneapi/dal/detail/profiler.hpp"
#include <iostream>

namespace oneapi::dal::backend {

Expand Down Expand Up @@ -175,6 +177,7 @@ inline unique_usm_ptr<void> make_unique_usm_shared(const sycl::queue& q, std::si
}

inline unique_usm_ptr<void> make_unique_usm_host(const sycl::queue& q, std::size_t size) {
ONEDAL_PROFILER_TASK(make_unique_usm_host);
return unique_usm_ptr<void>{ malloc_host(q, size), usm_deleter<void>{ q } };
}

Expand All @@ -197,6 +200,7 @@ inline unique_usm_ptr<T> make_unique_usm_shared(const sycl::queue& q, std::int64

template <typename T>
inline unique_usm_ptr<T> make_unique_usm_host(const sycl::queue& q, std::int64_t count) {
ONEDAL_PROFILER_TASK(make_unique_usm_host_with_template);
return unique_usm_ptr<T>{ malloc_host<T>(q, count), usm_deleter<T>{ q } };
}

Expand All @@ -220,12 +224,9 @@ inline sycl::event memcpy_host2usm(sycl::queue& queue,
std::size_t size,
const event_vector& deps = {}) {
ONEDAL_ASSERT(is_known_usm(queue, dest_usm));

// TODO: Remove additional copy to host usm memory once
// bug in `copy` with the host memory is fixed
auto tmp_usm_host = make_unique_usm_host(queue, size);
memcpy(tmp_usm_host.get(), src_host, size);
memcpy(queue, dest_usm, tmp_usm_host.get(), size, deps).wait_and_throw();
ONEDAL_PROFILER_TASK(memcpy_host2usm, queue);
std::cout<<"memcpy_host2usm"<<std::endl;
queue.memcpy(dest_usm, src_host, size, deps).wait_and_throw();
return {};
}

Expand All @@ -235,12 +236,9 @@ inline sycl::event memcpy_usm2host(sycl::queue& queue,
std::size_t size,
const event_vector& deps = {}) {
ONEDAL_ASSERT(is_known_usm(queue, src_usm));

// TODO: Remove additional copy to host usm memory once
// bug in `copy` with the host memory is fixed
auto tmp_usm_host = make_unique_usm_host(queue, size);
memcpy(queue, tmp_usm_host.get(), src_usm, size, deps).wait_and_throw();
memcpy(dest_host, tmp_usm_host.get(), size);
std::cout<<"memcpy_usm2host"<<std::endl;
ONEDAL_PROFILER_TASK(memcpy_usm2host, queue);
queue.memcpy(dest_host, src_usm, size, deps).wait_and_throw();
return {};
}

Expand Down Expand Up @@ -291,6 +289,7 @@ inline sycl::event copy_host2usm(sycl::queue& queue,
ONEDAL_ASSERT(count > 0);
const std::size_t n = detail::integral_cast<std::size_t>(count);
ONEDAL_ASSERT_MUL_OVERFLOW(std::size_t, sizeof(T), n);
std::cout<<"failed 3"<<std::endl;
return memcpy_host2usm(queue, dest_usm, src_host, sizeof(T) * n, deps);
}

Expand Down Expand Up @@ -386,9 +385,11 @@ inline sycl::event copy_all2all(sycl::queue& queue,
event = memcpy(queue, dest, src, sizeof(T) * n, deps);
}
else if (src_device_friendly) {
std::cout<<"failed 1"<<std::endl;
event = memcpy_usm2host(queue, dest, src, sizeof(T) * n, deps);
}
else if (dst_device_friendly) {
std::cout<<"failed 2"<<std::endl;
event = memcpy_host2usm(queue, dest, src, sizeof(T) * n, deps);
}
else {
Expand Down
Loading
Loading