Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Reduce the test scope for sort.cpp #670

Merged
merged 4 commits into from
Dec 24, 2021
Merged

Reduce the test scope for sort.cpp #670

Changes from 3 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
8a9bead
Reduce the test scope for sort.cpp
bader Dec 23, 2021
61a138b
Reduce work-group size and apply minor style fixes.
bader Dec 23, 2021
d2abf2a
Update SYCL/GroupAlgorithm/SYCL2020/sort.cpp
bader Dec 24, 2021
658c11e
Apply code review comments: use namespace alias instead of importing.
bader Dec 24, 2021
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
83 changes: 39 additions & 44 deletions SYCL/GroupAlgorithm/SYCL2020/sort.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,18 +2,16 @@
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out
//
// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t13.out

#include "support.h"
#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

#include <algorithm>
#include <iostream>
#include <random>
#include <vector>

namespace my_sycl = sycl::ext::oneapi::experimental;
using namespace sycl::ext::oneapi::experimental;

auto async_handler_ = [](sycl::exception_list ex_list) {
for (auto &ex : ex_list) {
Expand All @@ -38,14 +36,11 @@ struct CustomFunctor {
}
};

// we need it since using std::abs leads to compilation error
template <typename T> T my_abs(T x) { return x >= 0 ? x : -x; }
Copy link

@andreyfe1 andreyfe1 Dec 24, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm ok with the change if there are no compilation errors related to abs as it's mentioned in the comment

Copy link
Author

@bader bader Dec 24, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

As you can see in pre-commit logs, there no compilation errors. BTW, the comment is about std::abs, not sycl::abs.
I trusted your comment and didn't try to use std::abs here.

Copy link

@andreyfe1 andreyfe1 Dec 24, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see. Sorry, missed that there is sycl::abs now.

Copy link
Author

@bader bader Dec 24, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I checked std::abs and I see the problem.

llvm-test-suite/SYCL/GroupAlgorithm/SYCL2020/sort.cpp:40:10: error: call to 'abs' is ambiguous
  return std::abs(lhs - rhs) > epsilon;
         ^~~~~~~~
/llvm-test-suite/SYCL/GroupAlgorithm/SYCL2020/sort.cpp:49:9: note: in instantiation of function template specialization 'check<unsigned int>' requested here
    if (check(expected[i], got[i], epsilon)) {
        ^

Standard library doesn't provide abs overloads for unsigned integer types. It make sense, so the question is: why SYCL spec defines abs for unsigned integer types?

I suppose your code might not work as expected for unsigned types. I suggest you to think about it.

Copy link

@andreyfe1 andreyfe1 Dec 24, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for pointing it out. What is your opinion whether sycl::abs for unsigned integer types is a bug or a feature of SYCL?

Copy link
Author

@bader bader Dec 24, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Technically, I don't any problems with defining such function, but the use case is unclear. This function does nothing, but return its only argument.
I suppose it wasn't intention to define such function. It's just the side effect of listing all supported OpenCL types as allowed parameters for integer functions.
@intel/dpcpp-specification-reviewers, FYI.

Copy link

@gmlueck gmlueck Jan 4, 2022

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I captured this question as a Kronos issue against the SYCL spec: https://gitlab.khronos.org/sycl/Specification/-/issues/584


template <typename T> bool check(T lhs, T rhs, float epsilon) {
return my_abs(lhs - rhs) > epsilon;
return sycl::abs(lhs - rhs) > epsilon;
}
bool check(CustomType lhs, CustomType rhs, float epsilon) {
return my_abs(lhs.x - rhs.x) > epsilon;
return sycl::abs(lhs.x - rhs.x) > epsilon;
}

template <typename T>
Expand Down Expand Up @@ -105,7 +100,7 @@ int test_sort_over_group(sycl::queue &q, std::size_t local,

sycl::range<dim> local_range = get_range<dim>(local);

std::size_t local_memory_size = my_sycl::default_sorter<>::memory_required<T>(
std::size_t local_memory_size = default_sorter<>::memory_required<T>(
sycl::memory_scope::work_group, local_range);

if (local_memory_size >
Expand All @@ -129,24 +124,24 @@ int test_sort_over_group(sycl::queue &q, std::size_t local,
case 0:
if constexpr (std::is_same_v<Compare, std::less<T>> &&
!std::is_same_v<T, CustomType>)
aI1[local_id] = my_sycl::sort_over_group(
my_sycl::group_with_scratchpad(
aI1[local_id] = sort_over_group(
group_with_scratchpad(
id.get_group(),
sycl::span{&scratch[0], local_memory_size}),
aI1[local_id]);
break;
case 1:
aI1[local_id] = my_sycl::sort_over_group(
my_sycl::group_with_scratchpad(
aI1[local_id] = sort_over_group(
group_with_scratchpad(
id.get_group(),
sycl::span{&scratch[0], local_memory_size}),
aI1[local_id], comp);
break;
case 2:
aI1[local_id] = my_sycl::sort_over_group(
id.get_group(), aI1[local_id],
my_sycl::default_sorter<Compare>(
sycl::span{&scratch[0], local_memory_size}));
aI1[local_id] =
sort_over_group(id.get_group(), aI1[local_id],
default_sorter<Compare>(sycl::span{
&scratch[0], local_memory_size}));
break;
}
});
Expand All @@ -160,8 +155,8 @@ int test_joint_sort(sycl::queue &q, std::size_t n_items, std::size_t local,
auto n = bufI1.size();
auto n_groups = (n - 1) / n_items + 1;

std::size_t local_memory_size = my_sycl::default_sorter<>::memory_required<T>(
sycl::memory_scope::work_group, n);
std::size_t local_memory_size =
default_sorter<>::memory_required<T>(sycl::memory_scope::work_group, n);
if (local_memory_size >
q.get_device().template get_info<sycl::info::device::local_mem_size>())
std::cout << "local_memory_size = " << local_memory_size << ", available = "
Expand All @@ -187,27 +182,26 @@ int test_joint_sort(sycl::queue &q, std::size_t n_items, std::size_t local,
case 0:
if constexpr (std::is_same_v<Compare, std::less<T>> &&
!std::is_same_v<T, CustomType>)
my_sycl::joint_sort(
my_sycl::group_with_scratchpad(
id.get_group(),
sycl::span{&scratch[0], local_memory_size}),
ptr_keys,
ptr_keys + sycl::min(n_items, n - group_id * n_items));
joint_sort(group_with_scratchpad(
id.get_group(),
sycl::span{&scratch[0], local_memory_size}),
ptr_keys,
ptr_keys +
sycl::min(n_items, n - group_id * n_items));
break;
case 1:
my_sycl::joint_sort(
my_sycl::group_with_scratchpad(
id.get_group(),
sycl::span{&scratch[0], local_memory_size}),
ptr_keys,
ptr_keys + sycl::min(n_items, n - group_id * n_items), comp);
joint_sort(group_with_scratchpad(
id.get_group(),
sycl::span{&scratch[0], local_memory_size}),
ptr_keys,
ptr_keys + sycl::min(n_items, n - group_id * n_items),
comp);
break;
case 2:
my_sycl::joint_sort(
id.get_group(), ptr_keys,
ptr_keys + sycl::min(n_items, n - group_id * n_items),
my_sycl::default_sorter<Compare>(
sycl::span{&scratch[0], local_memory_size}));
joint_sort(id.get_group(), ptr_keys,
ptr_keys + sycl::min(n_items, n - group_id * n_items),
default_sorter<Compare>(
sycl::span{&scratch[0], local_memory_size}));
break;
}
});
Expand All @@ -217,7 +211,7 @@ int test_joint_sort(sycl::queue &q, std::size_t n_items, std::size_t local,

template <typename T, typename Compare>
int test_custom_sorter(sycl::queue &q, sycl::buffer<T> &bufI1, Compare comp) {
std::size_t local = 256;
std::size_t local = 4;
auto n = bufI1.size();
if (n > local)
return -1;
Expand All @@ -230,9 +224,8 @@ int test_custom_sorter(sycl::queue &q, sycl::buffer<T> &bufI1, Compare comp) {
sycl::nd_range<2>({local, 1}, {local, 1}), [=](sycl::nd_item<2> id) {
auto ptr = aI1.get_pointer();

my_sycl::joint_sort(
id.get_group(), ptr, ptr + n,
bubble_sorter<Compare>{comp, id.get_local_linear_id()});
joint_sort(id.get_group(), ptr, ptr + n,
bubble_sorter<Compare>{comp, id.get_local_linear_id()});
});
}).wait_and_throw();
return 1;
Expand All @@ -243,9 +236,11 @@ void run_sort(sycl::queue &q, std::vector<T> &in, std::size_t size,
Compare comp, int test_case, int sort_case) {
std::vector<T> in2(in.begin(), in.begin() + size);
std::vector<T> expected(in.begin(), in.begin() + size);
std::size_t local =
constexpr size_t work_size_limit = 4;
std::size_t local = std::min(
work_size_limit,
q.get_device()
.template get_info<sycl::info::device::max_work_group_size>();
.template get_info<sycl::info::device::max_work_group_size>());
local = std::min(local, size);
auto n_items = items_per_work_item * local;

Expand Down Expand Up @@ -354,7 +349,7 @@ int main(int argc, char *argv[]) {
return 0;
}

std::vector<int> sizes{1, 2, 64, 256, 1024, 2048, 4096};
std::vector<int> sizes{1, 12, 32};

for (int i = 0; i < sizes.size(); ++i) {
test_sort_by_type<std::int8_t>(q, sizes[i]);
Expand Down