Skip to content

[SYCL] Fix performance regression in parallel_for with using id class #5385

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

Merged
bader merged 4 commits into from
Feb 8, 2022
Merged

[SYCL] Fix performance regression in parallel_for with using id class #5385

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
78f2888
Fix performance regression in parallel_for with using id class
aobolensk Jan 25, 2022
a160066
Use _t and _v versions of type traits
aobolensk Jan 26, 2022
0ddb352
Add test for sycl::id type
aobolensk Jan 26, 2022
9a841b1
Add is_same_v and is_convertible_v to detail namespace
aobolensk Feb 7, 2022
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
13 changes: 8 additions & 5 deletions sycl/include/CL/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -917,11 +917,14 @@ class __SYCL_EXPORT handler {
}

template <int Dims, typename LambdaArgType> struct TransformUserItemType {
using type = typename std::conditional<
std::is_convertible<nd_item<Dims>, LambdaArgType>::value, nd_item<Dims>,
typename std::conditional<
std::is_convertible<item<Dims>, LambdaArgType>::value, item<Dims>,
LambdaArgType>::type>::type;
using type = typename std::conditional_t<
std::is_same<id<Dims>, LambdaArgType>::value, LambdaArgType,
typename std::conditional_t<
std::is_convertible<nd_item<Dims>, LambdaArgType>::value,
nd_item<Dims>,
typename std::conditional_t<
std::is_convertible<item<Dims>, LambdaArgType>::value,
item<Dims>, LambdaArgType>>>;
Copy link
Contributor

@keryell keryell Jan 27, 2022

Choose a reason for hiding this comment

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

We are using at least C++17, right?

Suggested change
using type = typename std::conditional_t<
std::is_same<id<Dims>, LambdaArgType>::value, LambdaArgType,
typename std::conditional_t<
std::is_convertible<nd_item<Dims>, LambdaArgType>::value,
nd_item<Dims>,
typename std::conditional_t<
std::is_convertible<item<Dims>, LambdaArgType>::value,
item<Dims>, LambdaArgType>>>;
using type = typename std::conditional_t<
std::is_same_v<id<Dims>, LambdaArgType>, LambdaArgType,
typename std::conditional_t<
std::is_convertible_v<nd_item<Dims>, LambdaArgType>,
nd_item<Dims>,
typename std::conditional_t<
std::is_convertible_v<item<Dims>, LambdaArgType>,
item<Dims>, LambdaArgType>>>;

Perhaps you can try whether you can remove the typename and see whether it still compiles.

Copy link
Contributor Author

@aobolensk aobolensk Jan 27, 2022

Choose a reason for hiding this comment

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

I have tried to replace std::is_same with std::is_same_v too, but it leads to basic_tests/stdcpp_compat.cpp failure (https://github.com/intel/llvm/runs/4948816096?check_suite_focus=true)

Copy link
Contributor

@keryell keryell Jan 28, 2022

Choose a reason for hiding this comment

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

By looking at the compilation line in your CI test I can see: "/__w/llvm/llvm/build/bin/clang" "--driver-mode=g++" "-Xsycl-target-backend=amdgcn-amd-amdhsa" "--offload-arch=gfx906" "-D_GLIBCXX_ASSERTIONS=1" "-std=c++14"
C++14 whaaaat? :-)
@bader do you still have some old C++14 while SYCL 2020 is requiring at least C++17?

Copy link
Contributor

@bader bader Jan 31, 2022

Choose a reason for hiding this comment

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

Yes.
https://github.com/intel/llvm/blob/sycl/sycl/test/basic_tests/stdcpp_compat.cpp#L13 - as you can see this comment says that some features can be unavailable. We are trying to transition old codes written in SYCL-1.2.1 softly.
New code requiring c++17 features can be added under #ifdef __cplusplus >= 201703L.

Copy link
Contributor

@keryell keryell Jan 31, 2022

Choose a reason for hiding this comment

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

Perhaps this is a good opportunity to extend https://github.com/intel/llvm/blob/54ede407edeb93b7e9334d1e725f48bf981b2965/sycl/include/CL/sycl/detail/stl_type_traits.hpp with a few _v ones in the meantime?

Copy link
Contributor Author

@aobolensk aobolensk Feb 7, 2022 •
edited
Loading

Choose a reason for hiding this comment

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

@bader do we need to add missing _v functions to stl_type_traits.hpp or we are going to leave it as it is now?

Copy link
Contributor

@bader bader Feb 7, 2022

Choose a reason for hiding this comment

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

do we need to add missing _v functions to stl_type_traits.hpp

@aobolensk, I think this is what @keryell requests. Do you have any objections?

Copy link
Contributor Author

@aobolensk aobolensk Feb 7, 2022

Choose a reason for hiding this comment

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

do we need to add missing _v functions to stl_type_traits.hpp

@aobolensk, I think this is what @keryell requests. Do you have any objections?

@bader No objections. I will add it.

};

/// Defines and invokes a SYCL kernel function for the specified range.
Expand Down
30 changes: 30 additions & 0 deletions sycl/test/basic_tests/parallel_for_type_check.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,30 @@
// RUN: %clangxx -fsycl -fsycl-device-only -D__SYCL_INTERNAL_API -O0 -c -emit-llvm -S -o - %s | FileCheck %s

// This test performs basic type check for sycl::id that is used in result type.

#include <CL/sycl.hpp>
#include <iostream>

int main() {
sycl::queue q;

// Initialize data array
const int sz = 16;
int data[sz] = {0};
for (int i = 0; i < sz; ++i) {
data[i] = i;
}

// Check user defined sycl::item wrapper
sycl::buffer<int> data_buf(data, sz);
q.submit([&](sycl::handler &h) {
auto buf_acc = data_buf.get_access<sycl::access::mode::read_write>(h);
h.parallel_for(
sycl::range<1>{sz},
// CHECK: cl{{.*}}sycl{{.*}}detail{{.*}}RoundedRangeKernel{{.*}}id{{.*}}main{{.*}}handler
[=](sycl::id<1> item) { buf_acc[item] += 1; });
});
q.wait();

return 0;
}