From 4fade17e4cdc58c36c76a8ed2f7a8078a492a37c Mon Sep 17 00:00:00 2001 From: "richard.north.iii" Date: Mon, 3 Feb 2025 11:32:30 -0800 Subject: [PATCH 01/26] Adding correlation distance metric --- .../backend/primitives/distance/metrics.hpp | 40 +++++++++++++++++++ 1 file changed, 40 insertions(+) diff --git a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp index b501c982e20..ef123405b79 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp +++ b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp @@ -116,4 +116,44 @@ struct chebyshev_metric : public metric_base { } }; +template +struct correlation_metric : public metric_base { +public: + correlation_metric(){} + template + Float operator()(InputIt1 first1, InputIt1 last1, InputIt2 first2) const { + constexpr Float zero = 0; + constexpr Float one = 1; + Float ip_acc = zero; + Float n1_acc = zero; + Float n2_acc = zero; + Float n1_sum = zero; + Float n1_sum = zero; + Float count = zero; + for (auto it1 = first1, it2 = first2; it1 != last1; ++it1, it2++) { + n1_sum += *it1; + n2_sum += *it2; + ++count; + } + + if (count == zero) + return Float(zero); + + const Float n1_mean = n1_num / count; + const Float n2_mean = n2_num / count; + + for (auto it1 = first1, it2 = first2; it1 != last1; ++it1, ++it2) { + const Float v1 = *it1 - n1_mean; + const Float v2 = *it2 - n2_mean; + n1_acc += (v1 * v1); + n2_acc += (v2 * v2); + ip_acc += (v1 * v2); + + } + const Float rsqn1 = one / std::sqrt(n1_acc); + const Float rsqn2 = one / std::sqrt(n2_acc); + return one - ip_acc * rsqn1 * rsqn2; + } +}; + } // namespace oneapi::dal::backend::primitives From bb6172107210ef130556c9172c76f85cda116d81 Mon Sep 17 00:00:00 2001 From: richardnorth3 Date: Mon, 3 Feb 2025 14:30:58 -0600 Subject: [PATCH 02/26] Update metrics.hpp --- cpp/oneapi/dal/backend/primitives/distance/metrics.hpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp index ef123405b79..5b0ceccea59 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp +++ b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp @@ -139,8 +139,8 @@ struct correlation_metric : public metric_base { if (count == zero) return Float(zero); - const Float n1_mean = n1_num / count; - const Float n2_mean = n2_num / count; + const Float n1_mean = n1_sum / count; + const Float n2_mean = n2_sum / count; for (auto it1 = first1, it2 = first2; it1 != last1; ++it1, ++it2) { const Float v1 = *it1 - n1_mean; @@ -148,7 +148,6 @@ struct correlation_metric : public metric_base { n1_acc += (v1 * v1); n2_acc += (v2 * v2); ip_acc += (v1 * v2); - } const Float rsqn1 = one / std::sqrt(n1_acc); const Float rsqn2 = one / std::sqrt(n2_acc); From cd7c506ef5199d61a5881174bc82490ef2955d3f Mon Sep 17 00:00:00 2001 From: North Iii Date: Fri, 14 Feb 2025 02:54:51 -0600 Subject: [PATCH 03/26] Update metrics.hpp with correlation_metric for GPU Signed-off-by: North Iii --- .../backend/primitives/distance/metrics.hpp | 60 +++++++++++++++++++ 1 file changed, 60 insertions(+) diff --git a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp index 5b0ceccea59..3178fe8bb9f 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp +++ b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp @@ -153,6 +153,66 @@ struct correlation_metric : public metric_base { const Float rsqn2 = one / std::sqrt(n2_acc); return one - ip_acc * rsqn1 * rsqn2; } + + template + sycl::event operator()(sycl::queue& q, + const ndview& u, + const ndview& v, + ndview& out, + const event_vector& deps = {}) const { + const std::int64_t n = u.get_dimension(0); + auto u_sum = ndarray::empty({ 1 }); + auto v_sum = ndarray::empty({ 1 }); + auto u_mean = ndarray::empty({ 1 }); + auto v_mean = ndarray::empty({ 1 }); + sycl::event evt1 = reduce_by_columns(q, u, u_sum, {}, {}, deps, true); + sycl::event evt2 = reduce_by_columns(q, v, v_sum, {}, {}, { evt1 }, true); + sycl::event evt3 = means(q, n, u_sum, u_mean, { evt2 }); + sycl::event evt4 = means(q, n, v_sum, v_mean, { evt3 }); + auto temp = ndarray::empty({ 3 }); + q.fill(temp, Float(0)).wait(); + sycl::event evt5 = q.submit([&](sycl::handler& h) { + h.depends_on({ evt4 }); + h.parallel_for(sycl::range<1>(n), [=](sycl::id<1> idx) { + const std::int64_t i = idx[0]; + const Float x = u.at(i, 0); + const Float y = v.at(i, 0); + const Float mu_x = u_mean.at(0); + const Float mu_y = v_mean.at(0); + const Float d1 = x - mu_x; + const Float d2 = y - mu_y; + sycl::atomic_ref + atomic_dot(temp.get_mutable_data()[0]); + sycl::atomic_ref + atomic_norm1(temp.get_mutable_data()[1]); + sycl::atomic_ref + atomic_norm2(temp.get_mutable_data()[2]); + atomic_dot.fetch_add(d1 * d2); + atomic_norm1.fetch_add(d1 * d1); + atomic_norm2.fetch_add(d2 * d2); + }); + }); + evt5.wait_and_throw(); + std::array host_temp; + q.memcpy(host_temp.data(), temp.get_mutable_data(), 3 * sizeof(Float)).wait(); + const Float dot = host_temp[0]; + const Float norm1 = host_temp[1]; + const Float norm2 = host_temp[2]; + const Float corr = + (norm1 > 0 && norm2 > 0) ? dot / (std::sqrt(norm1) * std::sqrt(norm2)) : Float(0); + const Float distance = 1 - corr; + out.get_mutable_data()[0] = distance; + return evt5; + } }; } // namespace oneapi::dal::backend::primitives From 20396751a2bd0bcce8fd3b8891d5f0543a5173c2 Mon Sep 17 00:00:00 2001 From: North Iii Date: Thu, 20 Feb 2025 09:51:09 -0600 Subject: [PATCH 04/26] Add correlation distance test Signed-off-by: North Iii --- .../backend/primitives/distance/metrics.hpp | 11 +- .../test/correlation_distance_dpc.cpp | 120 ++++++++++++++++++ 2 files changed, 129 insertions(+), 2 deletions(-) create mode 100644 cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp diff --git a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp index 3178fe8bb9f..3063f707e82 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp +++ b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp @@ -17,8 +17,13 @@ #pragma once #include +#include +#include #include "oneapi/dal/backend/primitives/distance/distance.hpp" +#include "oneapi/dal/backend/primitives/ndarray.hpp" +#include "oneapi/dal/backend/primitives/reduction.hpp" +#include "oneapi/dal/backend/primitives/cov.hpp" namespace oneapi::dal::backend::primitives { @@ -116,6 +121,7 @@ struct chebyshev_metric : public metric_base { } }; +#ifdef ONEDAL_DATA_PARALLEL template struct correlation_metric : public metric_base { public: @@ -165,8 +171,8 @@ struct correlation_metric : public metric_base { auto v_sum = ndarray::empty({ 1 }); auto u_mean = ndarray::empty({ 1 }); auto v_mean = ndarray::empty({ 1 }); - sycl::event evt1 = reduce_by_columns(q, u, u_sum, {}, {}, deps, true); - sycl::event evt2 = reduce_by_columns(q, v, v_sum, {}, {}, { evt1 }, true); + sycl::event evt1 = reduce_by_rows(q, u, u_sum, {}, {}, deps, true); + sycl::event evt2 = reduce_by_rows(q, v, v_sum, {}, {}, { evt1 }, true); sycl::event evt3 = means(q, n, u_sum, u_mean, { evt2 }); sycl::event evt4 = means(q, n, v_sum, v_mean, { evt3 }); auto temp = ndarray::empty({ 3 }); @@ -214,5 +220,6 @@ struct correlation_metric : public metric_base { return evt5; } }; +#endif } // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp new file mode 100644 index 00000000000..2907331ca1e --- /dev/null +++ b/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp @@ -0,0 +1,120 @@ +/******************************************************************************* +* Copyright 2022 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 +#include +#include + +#include "oneapi/dal/test/engine/common.hpp" +#include "oneapi/dal/test/engine/fixtures.hpp" +#include "oneapi/dal/test/engine/dataframe.hpp" + +#include "oneapi/dal/table/row_accessor.hpp" + +#include "oneapi/dal/backend/primitives/distance/distance.hpp" + +namespace oneapi::dal::backend::primitives::test { + +namespace te = dal::test::engine; +namespace pr = oneapi::dal::backend::primitives; + +using distance_types = std::tuple; + +template +class correlation_distance_test_random : public te::float_algo_fixture { +public: + void generate() { + r_count1_ = GENERATE(17, 31); + r_count2_ = GENERATE(7, 29); + c_count_ = GENERATE(3, 13); + generate_input(); + } + + te::table_id get_homogen_table_id() const { + return te::table_id::homogen(); + } + + auto output() { + return ndarray::zeros(this->get_queue(), { r_count1_, r_count2_ }); + } + + void generate_input() { + const auto input1_dataframe = GENERATE_DATAFRAME( + te::dataframe_builder{ r_count1_, c_count_ }.fill_uniform(-0.2, 0.5)); + this->input_table1_ = input1_dataframe.get_table(this->get_homogen_table_id()); + const auto input2_dataframe = GENERATE_DATAFRAME( + te::dataframe_builder{ r_count2_, c_count_ }.fill_uniform(-0.5, 1.0)); + this->input_table2_ = input2_dataframe.get_table(this->get_homogen_table_id()); + } + + void groundtruth_check(const ndview& out, const Float atol = 1.e-3) { + for (std::int64_t i = 0; i < r_count1_; ++i) { + const auto inp_row1 = + row_accessor{ input_table1_ }.pull(this->get_queue(), { i, i + 1 }); + for (std::int64_t j = 0; j < r_count2_; ++j) { + const auto inp_row2 = + row_accessor{ input_table2_ }.pull(this->get_queue(), + { j, j + 1 }); + Float ip = 0.0, qn = 0.0, tn = 0.0; + for (std::int64_t k = 0; k < c_count_; ++k) { + const auto q = inp_row1[k]; + const auto t = inp_row2[k]; + qn += q * q; + tn += t * t; + ip += q * t; + } + const auto gtv = Float(1.0) - ip / (std::sqrt(qn) * std::sqrt(tn)); + + const auto val = *(out.get_data() + out.get_leading_stride() * i + j); + const auto diff = gtv - val; + CAPTURE(gtv, val, i, j, r_count1_, r_count2_, c_count_); + REQUIRE(-atol <= diff); + CAPTURE(gtv, val, i, j, r_count1_, r_count2_, c_count_); + REQUIRE(diff <= atol); + } + } + } + + void test_distance() { + auto input1_arr = row_accessor{ input_table1_ }.pull(this->get_queue()); + auto input2_arr = row_accessor{ input_table2_ }.pull(this->get_queue()); + auto input1 = ndview::wrap(input1_arr.get_data(), { r_count1_, c_count_ }); + auto input2 = ndview::wrap(input2_arr.get_data(), { r_count2_, c_count_ }); + auto [output, output_event] = this->output(); + correlation_distance distance(this->get_queue()); + auto distance_event = distance(input1, input2, output, { output_event }); + distance_event.wait_and_throw(); + groundtruth_check(output); + } + +private: + table input_table1_; + table input_table2_; + std::int64_t c_count_; + std::int64_t r_count1_; + std::int64_t r_count2_; +}; + +TEMPLATE_LIST_TEST_M(correlation_distance_test_random, + "Randomly filled correlation-distance computation", + "[correlation][distance][small]", + distance_types) { + SKIP_IF(this->not_float64_friendly()); + this->generate(); + this->test_distance(); +} + +} // namespace oneapi::dal::backend::primitives::test From b4fd66a98103b6926fe2e83605d659fc8a55865f Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 01:27:41 -0600 Subject: [PATCH 05/26] Add supporting files for correlation distance Signed-off-by: North Iii --- .../distance/correlation_distance_dpc.cpp | 130 ++++++++++++++++ .../distance/correlation_distance_misc.hpp | 57 +++++++ .../correlation_distance_misc_dpc.cpp | 145 ++++++++++++++++++ .../backend/primitives/distance/metrics.hpp | 62 -------- 4 files changed, 332 insertions(+), 62 deletions(-) create mode 100644 cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp create mode 100644 cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp create mode 100644 cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp new file mode 100644 index 00000000000..af482af9702 --- /dev/null +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -0,0 +1,130 @@ +/******************************************************************************* +* Copyright 2022 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/backend/primitives/distance/distance.hpp" +#include "oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp" + +#include "oneapi/dal/backend/primitives/blas.hpp" +#include "oneapi/dal/backend/primitives/reduction.hpp" + +namespace oneapi::dal::backend::primitives { + +template +template +auto distance>::get_inversed_norms(const ndview& inp, + const event_vector& deps) const + -> inv_norms_res_t { + return compute_inversed_l2_norms(q_, inp, deps); +} + +template +template +sycl::event distance>::operator()(const ndview& inp1, + const ndview& inp2, + ndview& out, + const ndview& inp1_norms, + const ndview& inp2_norms, + const event_vector& deps) const { + auto ip_event = compute_correlation_inner_product(q_, inp1, inp2, out, deps); + return finalize_correlation(q_, inp1_norms, inp2_norms, out, { ip_event }); +} + +template +template +sycl::event distance>::operator()(const ndview& inp1, + const ndview& inp2, + ndview& out, + const event_vector& deps) const { + const std::int64_t n = inp1.get_dimension(0); + const std::int64_t p = inp1.get_dimension(1); + auto inp1_sum = ndarray::empty(q, { n }); + auto inp2_sum = ndarray::empty(q, { n }); + auto inp1_mean = ndarray::empty(q, { n }); + auto inp2_mean = ndarray::empty(q, { n }); + sycl::event evt1 = reduce_by_rows(q, + inp1, + inp1_sum, + reduction::plus{}, + reduction::identity{}, + deps); + sycl::event evt2 = reduce_by_rows(q, + inp2, + inp2_sum, + reduction::plus{}, + reduction::identity{}, + { evt1 }); + sycl::event evt3 = means(q, p, inp1_sum, inp1_mean, { evt2 }); + sycl::event evt4 = means(q, p, inp2_sum, inp2_mean, { evt3 }); + auto centered_inp1 = ndarray::empty(q, { n, p }); + auto centered_inp2 = ndarray::empty(q, { n, p }); + sycl::event evt5 = q.submit([&](sycl::handler& h) { + h.depends_on({ evt4 }); + auto inp1_acc = inp1.get_access(h); + auto inp2_acc = inp2.get_access(h); + auto inp1_mean_acc = inp1_mean.get_access(h); + auto inp2_mean_acc = inp2_mean.get_access(h); + auto centered1_acc = centered_inp1.get_access(h); + auto centered2_acc = centered_inp2.get_access(h); + + h.parallel_for(sycl::range<2>(n, p), [=](sycl::id<2> idx) { + const std::int64_t row = idx[0]; + const std::int64_t col = idx[1]; + centered1_acc(row, col) = inp1_acc(row, col) - inp1_mean_acc[row]; + centered2_acc(row, col) = inp2_acc(row, col) - inp2_mean_acc[row]; + }); + }); + evt5.wait(); + + auto [inv_norms1_array, inv_norms1_event] = get_inversed_norms(centered_inp1, { evt5 }); + auto [inv_norms2_array, inv_norms2_event] = get_inversed_norms(centered_inp2, { evt5 }); + return this->operator()(inp1, + inp2, + out, + inv_norms1_array, + inv_norms2_array, + { inv_norms1_event, inv_norms2_event }); +} + +#define INSTANTIATE(F, A, B) \ + template sycl::event distance>::operator()(const ndview&, \ + const ndview&, \ + ndview&, \ + const ndview&, \ + const ndview&, \ + const event_vector&) const; \ + template sycl::event distance>::operator()(const ndview&, \ + const ndview&, \ + ndview&, \ + const event_vector&) const; + +#define INSTANTIATE_B(F, A) \ + INSTANTIATE(F, A, ndorder::c) \ + INSTANTIATE(F, A, ndorder::f) \ + template std::tuple, sycl::event> \ + distance>::get_inversed_norms(const ndview& inp, \ + const event_vector& deps) const; + +#define INSTANTIATE_F(F) \ + INSTANTIATE_B(F, ndorder::c) \ + INSTANTIATE_B(F, ndorder::f) \ + template class distance>; + +INSTANTIATE_F(float); +INSTANTIATE_F(double); + +#undef INSTANTIATE + +} // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp new file mode 100644 index 00000000000..b49f047328b --- /dev/null +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp @@ -0,0 +1,57 @@ +/******************************************************************************* +* Copyright 2022 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. +*******************************************************************************/ + +#pragma once + +#include "oneapi/dal/backend/primitives/common.hpp" +#include "oneapi/dal/backend/primitives/ndarray.hpp" + +#include "oneapi/dal/backend/primitives/distance/distance.hpp" + +namespace oneapi::dal::backend::primitives { + +#ifdef ONEDAL_DATA_PARALLEL + +template +sycl::event compute_inversed_l2_norms(sycl::queue& q, + const ndview& inp, + ndview& out, + const event_vector& deps = {}); + +template +std::tuple, sycl::event> compute_inversed_l2_norms( + sycl::queue& q, + const ndview& inp, + const event_vector& deps = {}, + const sycl::usm::alloc& alloc = sycl::usm::alloc::device); + +template +sycl::event compute_correlation_inner_product(sycl::queue& q, + const ndview& inp1, + const ndview& inp2, + ndview& out, + const event_vector& deps = {}); + +template +sycl::event finalize_correlation(sycl::queue& q, + const ndview& inp1, + const ndview& inp2, + ndview& out, + const event_vector& deps = {}); + +#endif + +} // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp new file mode 100644 index 00000000000..00113091c97 --- /dev/null +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp @@ -0,0 +1,145 @@ +/******************************************************************************* +* Copyright 2022 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/distance/cosine_distance_misc.hpp" +#include "oneapi/dal/backend/primitives/distance/squared_l2_distance_misc.hpp" + +#include "oneapi/dal/backend/primitives/blas.hpp" +#include "oneapi/dal/backend/primitives/reduction.hpp" + +namespace oneapi::dal::backend::primitives { + +template +inline sycl::event inverse_l2_norms(sycl::queue& q, + ndview& out, + const event_vector& deps) { + ONEDAL_PROFILER_TASK(distance.inverse_l2_norms, q); + + ONEDAL_ASSERT(out.has_mutable_data()); + return q.submit([&](sycl::handler& h) { + h.depends_on(deps); + const auto count = out.get_count(); + const auto range = make_range_1d(count); + auto* const ptr = out.get_mutable_data(); + h.parallel_for(range, [=](sycl::id<1> idx) { + auto& ref = ptr[idx]; + ref = sycl::rsqrt(ref); + }); + }); +} + +template +sycl::event compute_inversed_l2_norms(sycl::queue& q, + const ndview& inp, + ndview& out, + const event_vector& deps) { + ONEDAL_ASSERT(inp.has_data()); + ONEDAL_ASSERT(out.has_mutable_data()); + auto sq_event = compute_squared_l2_norms(q, inp, out, deps); + return inverse_l2_norms(q, out, { sq_event }); +} + +template +std::tuple, sycl::event> compute_inversed_l2_norms( + sycl::queue& q, + const ndview& inp, + const event_vector& deps, + const sycl::usm::alloc& alloc) { + const auto n_samples = inp.get_dimension(0); + auto res_array = ndarray::empty(q, { n_samples }, alloc); + return { res_array, compute_inversed_l2_norms(q, inp, res_array, deps) }; +} + +template +sycl::event finalize_cosine(sycl::queue& q, + const ndview& inp1, + const ndview& inp2, + ndview& out, + const event_vector& deps) { + ONEDAL_PROFILER_TASK(distance.finalize_cosine, q); + + ONEDAL_ASSERT(inp1.has_data()); + ONEDAL_ASSERT(inp2.has_data()); + ONEDAL_ASSERT(out.has_mutable_data()); + const auto out_stride = out.get_leading_stride(); + const auto n_samples1 = inp1.get_dimension(0); + const auto n_samples2 = inp2.get_dimension(0); + ONEDAL_ASSERT(n_samples1 <= out.get_dimension(0)); + ONEDAL_ASSERT(n_samples2 <= out.get_dimension(1)); + const auto* const inp1_ptr = inp1.get_data(); + const auto* const inp2_ptr = inp2.get_data(); + auto* const out_ptr = out.get_mutable_data(); + const auto out_range = make_range_2d(n_samples1, n_samples2); + return q.submit([&](sycl::handler& h) { + h.depends_on(deps); + h.parallel_for(out_range, [=](sycl::id<2> idx) { + constexpr Float one = 1; + auto& out = *(out_ptr + out_stride * idx[0] + idx[1]); + out = one - out * inp1_ptr[idx[0]] * inp2_ptr[idx[1]]; + }); + }); +} + +template +sycl::event compute_cosine_inner_product(sycl::queue& q, + const ndview& inp1, + const ndview& inp2, + ndview& out, + const event_vector& deps) { + check_inputs(inp1, inp2, out); + auto event = gemm(q, inp1, inp2.t(), out, Float(+1.0), Float(0.0), deps); + // Workaround for abort in async mode. Should be removed later. + event.wait_and_throw(); + return event; +} + +#define INSTANTIATE(F, A, B) \ + template sycl::event compute_cosine_inner_product(sycl::queue&, \ + const ndview&, \ + const ndview&, \ + ndview&, \ + const event_vector&); + +#define INSTANTIATE_A(F, B) \ + INSTANTIATE(F, ndorder::c, B) \ + INSTANTIATE(F, ndorder::f, B) \ + template sycl::event compute_inversed_l2_norms(sycl::queue&, \ + const ndview&, \ + ndview&, \ + const event_vector&); \ + template std::tuple, sycl::event> compute_inversed_l2_norms( \ + sycl::queue&, \ + const ndview&, \ + const event_vector&, \ + const sycl::usm::alloc&); + +#define INSTANTIATE_F(F) \ + INSTANTIATE_A(F, ndorder::c) \ + INSTANTIATE_A(F, ndorder::f) \ + template sycl::event finalize_cosine(sycl::queue & q, \ + const ndview&, \ + const ndview&, \ + ndview&, \ + const event_vector&); + +INSTANTIATE_F(float); +INSTANTIATE_F(double); + +#undef INSTANTIATE + +} // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp index 3063f707e82..3e6d9c6cb70 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp +++ b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp @@ -121,7 +121,6 @@ struct chebyshev_metric : public metric_base { } }; -#ifdef ONEDAL_DATA_PARALLEL template struct correlation_metric : public metric_base { public: @@ -159,67 +158,6 @@ struct correlation_metric : public metric_base { const Float rsqn2 = one / std::sqrt(n2_acc); return one - ip_acc * rsqn1 * rsqn2; } - - template - sycl::event operator()(sycl::queue& q, - const ndview& u, - const ndview& v, - ndview& out, - const event_vector& deps = {}) const { - const std::int64_t n = u.get_dimension(0); - auto u_sum = ndarray::empty({ 1 }); - auto v_sum = ndarray::empty({ 1 }); - auto u_mean = ndarray::empty({ 1 }); - auto v_mean = ndarray::empty({ 1 }); - sycl::event evt1 = reduce_by_rows(q, u, u_sum, {}, {}, deps, true); - sycl::event evt2 = reduce_by_rows(q, v, v_sum, {}, {}, { evt1 }, true); - sycl::event evt3 = means(q, n, u_sum, u_mean, { evt2 }); - sycl::event evt4 = means(q, n, v_sum, v_mean, { evt3 }); - auto temp = ndarray::empty({ 3 }); - q.fill(temp, Float(0)).wait(); - sycl::event evt5 = q.submit([&](sycl::handler& h) { - h.depends_on({ evt4 }); - h.parallel_for(sycl::range<1>(n), [=](sycl::id<1> idx) { - const std::int64_t i = idx[0]; - const Float x = u.at(i, 0); - const Float y = v.at(i, 0); - const Float mu_x = u_mean.at(0); - const Float mu_y = v_mean.at(0); - const Float d1 = x - mu_x; - const Float d2 = y - mu_y; - sycl::atomic_ref - atomic_dot(temp.get_mutable_data()[0]); - sycl::atomic_ref - atomic_norm1(temp.get_mutable_data()[1]); - sycl::atomic_ref - atomic_norm2(temp.get_mutable_data()[2]); - atomic_dot.fetch_add(d1 * d2); - atomic_norm1.fetch_add(d1 * d1); - atomic_norm2.fetch_add(d2 * d2); - }); - }); - evt5.wait_and_throw(); - std::array host_temp; - q.memcpy(host_temp.data(), temp.get_mutable_data(), 3 * sizeof(Float)).wait(); - const Float dot = host_temp[0]; - const Float norm1 = host_temp[1]; - const Float norm2 = host_temp[2]; - const Float corr = - (norm1 > 0 && norm2 > 0) ? dot / (std::sqrt(norm1) * std::sqrt(norm2)) : Float(0); - const Float distance = 1 - corr; - out.get_mutable_data()[0] = distance; - return evt5; - } }; -#endif } // namespace oneapi::dal::backend::primitives From 22f03a889a89318eb2a4c64f536a7265764eccdf Mon Sep 17 00:00:00 2001 From: richardnorth3 Date: Mon, 24 Feb 2025 01:32:41 -0600 Subject: [PATCH 06/26] Update metrics.hpp --- cpp/oneapi/dal/backend/primitives/distance/metrics.hpp | 5 ----- 1 file changed, 5 deletions(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp index 3e6d9c6cb70..5b0ceccea59 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp +++ b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp @@ -17,13 +17,8 @@ #pragma once #include -#include -#include #include "oneapi/dal/backend/primitives/distance/distance.hpp" -#include "oneapi/dal/backend/primitives/ndarray.hpp" -#include "oneapi/dal/backend/primitives/reduction.hpp" -#include "oneapi/dal/backend/primitives/cov.hpp" namespace oneapi::dal::backend::primitives { From 93a1c984ec8b5877c3b4a2537e27ce9211bf93e1 Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 01:54:27 -0600 Subject: [PATCH 07/26] Update distance.hpp with correlation prototypes Signed-off-by: North Iii --- .../backend/primitives/distance/distance.hpp | 33 +++++++++++++++++++ 1 file changed, 33 insertions(+) diff --git a/cpp/oneapi/dal/backend/primitives/distance/distance.hpp b/cpp/oneapi/dal/backend/primitives/distance/distance.hpp index f38668bab85..0c614114d96 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/distance.hpp +++ b/cpp/oneapi/dal/backend/primitives/distance/distance.hpp @@ -103,6 +103,36 @@ class distance> { sycl::queue& q_; }; +template +class distance> { +public: + distance(sycl::queue& q) : q_{ q } {}; + + template + sycl::event operator()(const ndview& inp1, + const ndview& inp2, + ndview& out, + const event_vector& deps = {}) const; + + template + sycl::event operator()(const ndview& inp1, + const ndview& inp2, + ndview& out, + const ndview& inp1_norms, + const ndview& inp2_norms, + const event_vector& deps = {}) const; + +protected: + using inv_norms_res_t = std::tuple, sycl::event>; + + template + inv_norms_res_t get_inversed_norms(const ndview& inp, + const event_vector& deps = {}) const; + +private: + sycl::queue& q_; +}; + template using lp_distance = distance>; @@ -115,6 +145,9 @@ using cosine_distance = distance>; template using chebyshev_distance = distance>; +template +using correlation_distance = distance>; + template void check_inputs(const ndview& inp1, const ndview& inp2, From a53164d256ee4355b2fbed5b985574084fb1e13c Mon Sep 17 00:00:00 2001 From: richardnorth3 Date: Mon, 24 Feb 2025 02:46:18 -0600 Subject: [PATCH 08/26] Update metrics.hpp --- cpp/oneapi/dal/backend/primitives/distance/metrics.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp index 5b0ceccea59..50bd121fc82 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp +++ b/cpp/oneapi/dal/backend/primitives/distance/metrics.hpp @@ -128,7 +128,7 @@ struct correlation_metric : public metric_base { Float n1_acc = zero; Float n2_acc = zero; Float n1_sum = zero; - Float n1_sum = zero; + Float n2_sum = zero; Float count = zero; for (auto it1 = first1, it2 = first2; it1 != last1; ++it1, it2++) { n1_sum += *it1; From 7ade1610e6ceb20800d79a15369ee9b6b8370af1 Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 03:03:16 -0600 Subject: [PATCH 09/26] Update correlation_distance_dpc.cpp Signed-off-by: North Iii --- .../distance/correlation_distance_dpc.cpp | 22 +++++++++---------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp index af482af9702..f319bcdaf7e 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -50,27 +50,27 @@ sycl::event distance>::operator()(const ndview< const event_vector& deps) const { const std::int64_t n = inp1.get_dimension(0); const std::int64_t p = inp1.get_dimension(1); - auto inp1_sum = ndarray::empty(q, { n }); - auto inp2_sum = ndarray::empty(q, { n }); - auto inp1_mean = ndarray::empty(q, { n }); - auto inp2_mean = ndarray::empty(q, { n }); - sycl::event evt1 = reduce_by_rows(q, + auto inp1_sum = ndarray::empty(q_, { n }); + auto inp2_sum = ndarray::empty(q_, { n }); + auto inp1_mean = ndarray::empty(q_, { n }); + auto inp2_mean = ndarray::empty(q_, { n }); + sycl::event evt1 = reduce_by_rows(q_, inp1, inp1_sum, reduction::plus{}, reduction::identity{}, deps); - sycl::event evt2 = reduce_by_rows(q, + sycl::event evt2 = reduce_by_rows(q_, inp2, inp2_sum, reduction::plus{}, reduction::identity{}, { evt1 }); - sycl::event evt3 = means(q, p, inp1_sum, inp1_mean, { evt2 }); - sycl::event evt4 = means(q, p, inp2_sum, inp2_mean, { evt3 }); - auto centered_inp1 = ndarray::empty(q, { n, p }); - auto centered_inp2 = ndarray::empty(q, { n, p }); - sycl::event evt5 = q.submit([&](sycl::handler& h) { + sycl::event evt3 = means(q_, p, inp1_sum, inp1_mean, { evt2 }); + sycl::event evt4 = means(q_, p, inp2_sum, inp2_mean, { evt3 }); + auto centered_inp1 = ndarray::empty(q_, { n, p }); + auto centered_inp2 = ndarray::empty(q_, { n, p }); + sycl::event evt5 = q_.submit([&](sycl::handler& h) { h.depends_on({ evt4 }); auto inp1_acc = inp1.get_access(h); auto inp2_acc = inp2.get_access(h); From c56b2097c5a8ae184ca01e5185dd90af7b835869 Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 03:08:55 -0600 Subject: [PATCH 10/26] Update correlation_distance_dpc.cpp Signed-off-by: North Iii --- .../distance/correlation_distance_dpc.cpp | 23 ++++++------------- 1 file changed, 7 insertions(+), 16 deletions(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp index f319bcdaf7e..70f10759e57 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -44,28 +44,19 @@ sycl::event distance>::operator()(const ndview< template template -sycl::event distance>::operator()(const ndview& inp1, - const ndview& inp2, - ndview& out, - const event_vector& deps) const { +sycl::event distance>::operator()( + const ndview& inp1, + const ndview& inp2, + ndview& out, + const event_vector& deps) const { const std::int64_t n = inp1.get_dimension(0); const std::int64_t p = inp1.get_dimension(1); auto inp1_sum = ndarray::empty(q_, { n }); auto inp2_sum = ndarray::empty(q_, { n }); auto inp1_mean = ndarray::empty(q_, { n }); auto inp2_mean = ndarray::empty(q_, { n }); - sycl::event evt1 = reduce_by_rows(q_, - inp1, - inp1_sum, - reduction::plus{}, - reduction::identity{}, - deps); - sycl::event evt2 = reduce_by_rows(q_, - inp2, - inp2_sum, - reduction::plus{}, - reduction::identity{}, - { evt1 }); + sycl::event evt1 = reduce_by_rows(q_, inp1, inp1_sum, {}, {}, deps); + sycl::event evt2 = reduce_by_rows(q_, inp2, inp2_sum, {}, {}, { evt1 }); sycl::event evt3 = means(q_, p, inp1_sum, inp1_mean, { evt2 }); sycl::event evt4 = means(q_, p, inp2_sum, inp2_mean, { evt3 }); auto centered_inp1 = ndarray::empty(q_, { n, p }); From 0b45447bbadbf7e810f4106f8d49df8c1ccb179b Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 03:12:06 -0600 Subject: [PATCH 11/26] Update correlation_distance_dpc.cpp Signed-off-by: North Iii --- .../backend/primitives/distance/correlation_distance_dpc.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp index 70f10759e57..8f819f43758 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -20,6 +20,8 @@ #include "oneapi/dal/backend/primitives/blas.hpp" #include "oneapi/dal/backend/primitives/reduction.hpp" +#include + namespace oneapi::dal::backend::primitives { template From 9b09822dd84b8eb2ee555e08b4576c8ce7d9c0e2 Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 03:16:56 -0600 Subject: [PATCH 12/26] Update correlation_distance_dpc.cpp Signed-off-by: North Iii --- .../backend/primitives/distance/correlation_distance_dpc.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp index 8f819f43758..ccdb62cf5e1 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -19,6 +19,9 @@ #include "oneapi/dal/backend/primitives/blas.hpp" #include "oneapi/dal/backend/primitives/reduction.hpp" +#include "oneapi/dal/backend/primitives/ndarray.hpp" +#include "oneapi/dal/backend/primitives/stat/cov.hpp" + #include From cc527925c91ab83593982f742e6186b6240e0958 Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 11:00:25 -0600 Subject: [PATCH 13/26] Update bazel BUILD file with cov deps Signed-off-by: North Iii --- cpp/oneapi/dal/backend/primitives/distance/BUILD | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/oneapi/dal/backend/primitives/distance/BUILD b/cpp/oneapi/dal/backend/primitives/distance/BUILD index 8f96e673eb0..8735f841e5b 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/BUILD +++ b/cpp/oneapi/dal/backend/primitives/distance/BUILD @@ -11,6 +11,7 @@ dal_module( "@onedal//cpp/oneapi/dal/backend/primitives:blas", "@onedal//cpp/oneapi/dal/backend/primitives:common", "@onedal//cpp/oneapi/dal/backend/primitives:reduction", + "@onedal//cpp/oneapi/dal/backend/primitives:cov", ], ) From 8232af440912ff8bd62da7257d6f6c2951ead17f Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 11:05:49 -0600 Subject: [PATCH 14/26] Update bazel BUILD file Signed-off-by: North Iii --- cpp/oneapi/dal/backend/primitives/distance/BUILD | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/BUILD b/cpp/oneapi/dal/backend/primitives/distance/BUILD index 8735f841e5b..d39785326c2 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/BUILD +++ b/cpp/oneapi/dal/backend/primitives/distance/BUILD @@ -11,7 +11,7 @@ dal_module( "@onedal//cpp/oneapi/dal/backend/primitives:blas", "@onedal//cpp/oneapi/dal/backend/primitives:common", "@onedal//cpp/oneapi/dal/backend/primitives:reduction", - "@onedal//cpp/oneapi/dal/backend/primitives:cov", + "@onedal//cpp/oneapi/dal/backend/primitives/stat:cov", ], ) From 8670a5af860e1c0113e663fb804c413c49a4fe29 Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 11:10:54 -0600 Subject: [PATCH 15/26] Update bazel BUILD file Signed-off-by: North Iii --- cpp/oneapi/dal/backend/primitives/distance/BUILD | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/BUILD b/cpp/oneapi/dal/backend/primitives/distance/BUILD index d39785326c2..613f7b41f08 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/BUILD +++ b/cpp/oneapi/dal/backend/primitives/distance/BUILD @@ -11,7 +11,7 @@ dal_module( "@onedal//cpp/oneapi/dal/backend/primitives:blas", "@onedal//cpp/oneapi/dal/backend/primitives:common", "@onedal//cpp/oneapi/dal/backend/primitives:reduction", - "@onedal//cpp/oneapi/dal/backend/primitives/stat:cov", + "@onedal//cpp/oneapi/dal/backend/primitives:stat", ], ) From 7d88579fcf94f5f2bcfba65f4b4441cf6de5ebc5 Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 11:15:57 -0600 Subject: [PATCH 16/26] Update correlation_distance_dpc.cpp Signed-off-by: North Iii --- .../backend/primitives/distance/correlation_distance_dpc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp index ccdb62cf5e1..36c155e6608 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -23,7 +23,7 @@ #include "oneapi/dal/backend/primitives/stat/cov.hpp" -#include +#include namespace oneapi::dal::backend::primitives { From 4740ad63f5da7ece22940c552252ad4a6d6659d3 Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 11:28:58 -0600 Subject: [PATCH 17/26] Update correlation_distance_dpc.cpp Signed-off-by: North Iii --- .../primitives/distance/correlation_distance_dpc.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp index 36c155e6608..7877fdb620e 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -68,12 +68,12 @@ sycl::event distance>::operator()( auto centered_inp2 = ndarray::empty(q_, { n, p }); sycl::event evt5 = q_.submit([&](sycl::handler& h) { h.depends_on({ evt4 }); - auto inp1_acc = inp1.get_access(h); - auto inp2_acc = inp2.get_access(h); - auto inp1_mean_acc = inp1_mean.get_access(h); - auto inp2_mean_acc = inp2_mean.get_access(h); - auto centered1_acc = centered_inp1.get_access(h); - auto centered2_acc = centered_inp2.get_access(h); + auto inp1_acc = inp1.get_data(); + auto inp2_acc = inp2.get_data(); + auto inp1_mean_acc = inp1_mean.get_data(); + auto inp2_mean_acc = inp2_mean.get_data(); + auto centered1_acc = centered_inp1.get_mutable_data(); + auto centered2_acc = centered_inp2.get_mutable_data(); h.parallel_for(sycl::range<2>(n, p), [=](sycl::id<2> idx) { const std::int64_t row = idx[0]; From 069a7870e0f7c54c8e8cc89c7a4fa3b7b8b4ebc4 Mon Sep 17 00:00:00 2001 From: North Iii Date: Mon, 24 Feb 2025 11:30:48 -0600 Subject: [PATCH 18/26] Update correlation_distance_dpc.cpp Signed-off-by: North Iii --- .../backend/primitives/distance/correlation_distance_dpc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp index 7877fdb620e..263ba783dd1 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -18,7 +18,7 @@ #include "oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp" #include "oneapi/dal/backend/primitives/blas.hpp" -#include "oneapi/dal/backend/primitives/reduction.hpp" +#include "oneapi/dal/backend/primitives/reduction/reduction.hpp" #include "oneapi/dal/backend/primitives/ndarray.hpp" #include "oneapi/dal/backend/primitives/stat/cov.hpp" From 325e01b1e9c92b17ebe44b0f77b7ca465d7f56f3 Mon Sep 17 00:00:00 2001 From: richardnorth3 Date: Tue, 25 Feb 2025 08:13:23 -0600 Subject: [PATCH 19/26] Update cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp Co-authored-by: Victoriya Fedotova --- .../backend/primitives/distance/correlation_distance_dpc.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp index 263ba783dd1..32f17127897 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -60,8 +60,8 @@ sycl::event distance>::operator()( auto inp2_sum = ndarray::empty(q_, { n }); auto inp1_mean = ndarray::empty(q_, { n }); auto inp2_mean = ndarray::empty(q_, { n }); - sycl::event evt1 = reduce_by_rows(q_, inp1, inp1_sum, {}, {}, deps); - sycl::event evt2 = reduce_by_rows(q_, inp2, inp2_sum, {}, {}, { evt1 }); + sycl::event evt1 = reduce_by_rows(q_, inp1, inp1_sum, sum{}, identity{}, deps); + sycl::event evt2 = reduce_by_rows(q_, inp2, inp2_sum, sum{}, identity{}, { evt1 }); sycl::event evt3 = means(q_, p, inp1_sum, inp1_mean, { evt2 }); sycl::event evt4 = means(q_, p, inp2_sum, inp2_mean, { evt3 }); auto centered_inp1 = ndarray::empty(q_, { n, p }); From 01c4b72324d84aba29797db009e65a8dd891a0af Mon Sep 17 00:00:00 2001 From: richardnorth3 Date: Tue, 25 Feb 2025 08:13:54 -0600 Subject: [PATCH 20/26] Update cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp Co-authored-by: Victoriya Fedotova --- .../backend/primitives/distance/correlation_distance_dpc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp index 32f17127897..f4bfbd88e7a 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2022 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. From f888e7af2458b1a51a1160d7b5c489859d3d73e6 Mon Sep 17 00:00:00 2001 From: "richard.north.iii" Date: Tue, 25 Feb 2025 07:05:50 -0800 Subject: [PATCH 21/26] Update correlation distance test and header Signed-off-by: richard.north.iii --- .../backend/primitives/distance/correlation_distance_dpc.cpp | 4 ++-- .../primitives/distance/test/correlation_distance_dpc.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp index f4bfbd88e7a..0e2f1c16b81 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_dpc.cpp @@ -78,8 +78,8 @@ sycl::event distance>::operator()( h.parallel_for(sycl::range<2>(n, p), [=](sycl::id<2> idx) { const std::int64_t row = idx[0]; const std::int64_t col = idx[1]; - centered1_acc(row, col) = inp1_acc(row, col) - inp1_mean_acc[row]; - centered2_acc(row, col) = inp2_acc(row, col) - inp2_mean_acc[row]; + centered1_acc[row * p + col] = inp1_acc[row * p + col] - inp1_mean_acc[row]; + centered2_acc[row * p + col] = inp2_acc[row * p + col] - inp2_mean_acc[row]; }); }); evt5.wait(); diff --git a/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp index 2907331ca1e..ab4d83957b9 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2022 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. From 71609fccfb9aa21b2acab22f48744d47294fe99c Mon Sep 17 00:00:00 2001 From: North Iii Date: Tue, 25 Feb 2025 09:09:52 -0600 Subject: [PATCH 22/26] Update correlation distance test Signed-off-by: North Iii --- .../test/correlation_distance_dpc.cpp | 20 +++++++++++++++---- 1 file changed, 16 insertions(+), 4 deletions(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp index ab4d83957b9..843b2ef815f 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp @@ -68,15 +68,27 @@ class correlation_distance_test_random : public te::float_algo_fixture { const auto inp_row2 = row_accessor{ input_table2_ }.pull(this->get_queue(), { j, j + 1 }); - Float ip = 0.0, qn = 0.0, tn = 0.0; + Float mean1 = 0.0, mean2 = 0.0; for (std::int64_t k = 0; k < c_count_; ++k) { - const auto q = inp_row1[k]; - const auto t = inp_row2[k]; + mean1 += inp_row1[k]; + mean2 += inp_row2[k]; + } + mean1 /= c_count_; + mean2 /= c_count_; + + Float numerator = 0.0, denom1 = 0.0, denom2 = 0.0; + for (std::int64_t k = 0; k < c_count_; ++k) { + const Float q = inp_row1[k] - mean1; + const Float t = inp_row2[k] - mean2; qn += q * q; tn += t * t; ip += q * t; } - const auto gtv = Float(1.0) - ip / (std::sqrt(qn) * std::sqrt(tn)); + + Float gtv = 1.0; + if (qn > 0 && tn > 0) { + gtv -= ip / (std::sqrt(qn) * std::sqrt(tn)); + } const auto val = *(out.get_data() + out.get_leading_stride() * i + j); const auto diff = gtv - val; From 8008137753adbdcf959d11c4a7f898b0bc5dc15c Mon Sep 17 00:00:00 2001 From: richardnorth3 Date: Tue, 25 Feb 2025 09:11:33 -0600 Subject: [PATCH 23/26] Update correlation_distance_misc.hpp --- .../backend/primitives/distance/correlation_distance_misc.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp index b49f047328b..15265716012 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2022 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. From dbbbc3e716597a5f19c7731ecc27db5463ee5182 Mon Sep 17 00:00:00 2001 From: richardnorth3 Date: Tue, 25 Feb 2025 09:11:59 -0600 Subject: [PATCH 24/26] Update correlation_distance_misc_dpc.cpp --- .../primitives/distance/correlation_distance_misc_dpc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp index 00113091c97..569a9246f1c 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2022 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. From a862a7458e4b4aa390dc9d9251aaa9f4a0afe776 Mon Sep 17 00:00:00 2001 From: "richard.north.iii" Date: Tue, 25 Feb 2025 07:14:41 -0800 Subject: [PATCH 25/26] Update correlation distance test Signed-off-by: richard.north.iii --- .../primitives/distance/test/correlation_distance_dpc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp index 843b2ef815f..5249ef3bac5 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/test/correlation_distance_dpc.cpp @@ -76,7 +76,7 @@ class correlation_distance_test_random : public te::float_algo_fixture { mean1 /= c_count_; mean2 /= c_count_; - Float numerator = 0.0, denom1 = 0.0, denom2 = 0.0; + Float ip = 0.0, qn = 0.0, tn = 0.0; for (std::int64_t k = 0; k < c_count_; ++k) { const Float q = inp_row1[k] - mean1; const Float t = inp_row2[k] - mean2; From 68e8c4f014ceb6469fb99546a308c4fa2b10ed91 Mon Sep 17 00:00:00 2001 From: North Iii Date: Tue, 25 Feb 2025 09:35:19 -0600 Subject: [PATCH 26/26] Update correlation_distance_misc_dpc.cpp Signed-off-by: North Iii --- .../distance/correlation_distance_misc_dpc.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp index 569a9246f1c..36a3e6e1800 100644 --- a/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/distance/correlation_distance_misc_dpc.cpp @@ -16,7 +16,7 @@ #include "oneapi/dal/detail/profiler.hpp" -#include "oneapi/dal/backend/primitives/distance/cosine_distance_misc.hpp" +#include "oneapi/dal/backend/primitives/distance/correlation_distance_misc.hpp" #include "oneapi/dal/backend/primitives/distance/squared_l2_distance_misc.hpp" #include "oneapi/dal/backend/primitives/blas.hpp" @@ -66,12 +66,12 @@ std::tuple, sycl::event> compute_inversed_l2_norms( } template -sycl::event finalize_cosine(sycl::queue& q, +sycl::event finalize_correlation(sycl::queue& q, const ndview& inp1, const ndview& inp2, ndview& out, const event_vector& deps) { - ONEDAL_PROFILER_TASK(distance.finalize_cosine, q); + ONEDAL_PROFILER_TASK(distance.finalize_correlation, q); ONEDAL_ASSERT(inp1.has_data()); ONEDAL_ASSERT(inp2.has_data()); @@ -96,7 +96,7 @@ sycl::event finalize_cosine(sycl::queue& q, } template -sycl::event compute_cosine_inner_product(sycl::queue& q, +sycl::event compute_correlation_inner_product(sycl::queue& q, const ndview& inp1, const ndview& inp2, ndview& out, @@ -109,7 +109,7 @@ sycl::event compute_cosine_inner_product(sycl::queue& q, } #define INSTANTIATE(F, A, B) \ - template sycl::event compute_cosine_inner_product(sycl::queue&, \ + template sycl::event compute_correlation_inner_product(sycl::queue&, \ const ndview&, \ const ndview&, \ ndview&, \ @@ -131,7 +131,7 @@ sycl::event compute_cosine_inner_product(sycl::queue& q, #define INSTANTIATE_F(F) \ INSTANTIATE_A(F, ndorder::c) \ INSTANTIATE_A(F, ndorder::f) \ - template sycl::event finalize_cosine(sycl::queue & q, \ + template sycl::event finalize_correlation(sycl::queue & q, \ const ndview&, \ const ndview&, \ ndview&, \