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

Fix "loop not vectorized" warning when running with newer icpx compiler #2818

Merged
merged 4 commits into from
Jun 24, 2024
Merged

Fix "loop not vectorized" warning when running with newer icpx compiler #2818

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
c110bd1
Fix dpc vectorization error when running with new compiler
Vika-F Jun 13, 2024
bbfff89
Update cpp/oneapi/dal/table/backend/csr_kernels.cpp
Vika-F Jun 14, 2024
51a3fc2
Fix nd_range size
Vika-F Jun 17, 2024
52e1c39
simplify sycl range size computations
Vika-F Jun 18, 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
14 changes: 11 additions & 3 deletions cpp/oneapi/dal/table/backend/csr_kernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,7 @@
* limitations under the License.
*******************************************************************************/

#include "oneapi/dal/backend/common.hpp"
#include "oneapi/dal/table/backend/csr_kernels.hpp"
#include "oneapi/dal/table/backend/convert.hpp"

Expand Down Expand Up @@ -411,6 +412,12 @@ bool is_sorted(sycl::queue& queue,

sycl::buffer<std::int64_t, 1> count_buf(&count_descending_pairs, sycl::range<1>(1));

const auto count_m1 = count - 1LL;
const auto wg_size = dal::backend::device_max_wg_size(queue);
const size_t count_m1_unsigned = static_cast<size_t>(count_m1);

const size_t wg_count = (count_m1 + wg_size - 1) / wg_size;

// count the number of pairs of the subsequent elements in the data array that are sorted
// in desccending order using sycl::reduction
queue
Expand All @@ -419,10 +426,11 @@ bool is_sorted(sycl::queue& queue,
auto count_descending_reduction =
sycl::reduction(count_buf, cgh, sycl::ext::oneapi::plus<std::int64_t>());

cgh.parallel_for(sycl::range<1>{ dal::detail::integral_cast<std::size_t>(count - 1) },
cgh.parallel_for(sycl::nd_range<1>{ wg_count * wg_size, wg_size },
count_descending_reduction,
[=](sycl::id<1> i, auto& count_descending) {
if (data[i] > data[i + 1])
[=](sycl::nd_item<1> idx, auto& count_descending) {
const auto i = idx.get_global_id(0);
if (i < count_m1_unsigned && data[i + 1] < data[i])
count_descending.combine(1);
});
})
Expand Down