Skip to content

[SYCL][InvokeSimd] Allow callables to return uniform #10714

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
merged 1 commit into from
Aug 8, 2023
Merged

[SYCL][InvokeSimd] Allow callables to return uniform #10714

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
1 change: 1 addition & 0 deletions llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,6 +63,7 @@ static const char *LegalSYCLFunctions[] = {
"^sycl::_V1::ext::oneapi::sub_group::.+",
"^sycl::_V1::ext::oneapi::experimental::spec_constant<.+>::.+",
"^sycl::_V1::ext::oneapi::experimental::this_sub_group",
"^sycl::_V1::ext::oneapi::experimental::uniform<.+>::.+",
"^sycl::_V1::ext::oneapi::bfloat16::.+",
"^sycl::_V1::ext::oneapi::experimental::if_architecture_is"};

Expand Down
3 changes: 2 additions & 1 deletion sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -371,7 +371,8 @@ constexpr bool has_struct_arg(Ret (*)(Args...)) {

template <typename Ret, typename... Args>
constexpr bool has_struct_ret(Ret (*)(Args...)) {
return std::is_class_v<Ret> && !is_simd_or_mask_type<Ret>::value;
return std::is_class_v<Ret> && !is_simd_or_mask_type<Ret>::value &&
!is_uniform_type<Ret>::value;
}

template <typename Ret, typename... Args>
Expand Down
30 changes: 27 additions & 3 deletions sycl/test-e2e/InvokeSimd/Spec/uniform_retval.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,12 @@
//
// VISALTO enable run
// RUN: env IGC_VISALTO=63 IGC_VCSaveStackCallLinkage=1 IGC_VCDirectCallsOnly=1 %{run} %t.out
//
// RUN: %{build} -DUNIFORM_RET_TYPE -fno-sycl-device-code-split-esimd -Xclang -fsycl-allow-func-ptr -o %t2.out
// RUN: env IGC_VCSaveStackCallLinkage=1 IGC_VCDirectCallsOnly=1 %{run} %t2.out
//
// VISALTO enable run
// RUN: env IGC_VISALTO=63 IGC_VCSaveStackCallLinkage=1 IGC_VCDirectCallsOnly=1 %{run} %t2.out

/*
* Test case #1
Expand Down Expand Up @@ -98,17 +104,35 @@ template <class T>
* returning the scalar as a SIMD type seems to work fine.
*/
template <class T>
__attribute__((always_inline)) T
__attribute__((always_inline))
#ifdef UNIFORM_RET_TYPE
uniform<T>
#else
T
#endif
ESIMD_CALLEE_return_uniform_scalar(esimd::simd<T, VL> x,
T n) SYCL_ESIMD_FUNCTION {
#ifdef UNIFORM_RET_TYPE
return uniform<T>{n};
#else
return n;
#endif
}

template <class T>
[[intel::device_indirectly_callable]] SYCL_EXTERNAL
T __regcall SIMD_CALLEE_return_uniform_scalar(simd<T, VL> x,
T n) SYCL_ESIMD_FUNCTION {
#ifdef UNIFORM_RET_TYPE
uniform<T>
#else
T
#endif
__regcall SIMD_CALLEE_return_uniform_scalar(simd<T, VL> x,
T n) SYCL_ESIMD_FUNCTION {
#ifdef UNIFORM_RET_TYPE
uniform<T> r = ESIMD_CALLEE_return_uniform_scalar<T>(x, n);
#else
T r = ESIMD_CALLEE_return_uniform_scalar<T>(x, n);
#endif
return r;
}

Expand Down
29 changes: 29 additions & 0 deletions sycl/test/invoke_simd/return-type-uniform.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,29 @@
// RUN: %clangxx -fsycl -fno-sycl-device-code-split-esimd -Xclang -fsycl-allow-func-ptr %s -o /dev/null
#include <sycl/ext/intel/esimd.hpp>
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>
#include <sycl/sycl.hpp>

using namespace sycl::ext::oneapi::experimental;
using namespace sycl;
namespace esimd = sycl::ext::intel::esimd;

[[intel::device_indirectly_callable]] uniform<int>
callee(simd<int, 8>) SYCL_ESIMD_FUNCTION {
return uniform<int>(5);
}
Comment on lines +10 to +13
Copy link
Contributor

@v-klochkov v-klochkov Aug 7, 2023 •
edited
Loading

Choose a reason for hiding this comment

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

Using 'uniform' in ESIMD context does not seem needed at all (IMO).

int/*uniform_return*/  callee(simd< int, 8 > vec, int uniform_param) SYCL_ESIMD_FUNCTION"

is enough to express the uniform-ness of the returned value and the second argument.

I suppose this patch is needed because of this invoke_simd SPEC statements that implicitly tells using uniform in ESIMD context is Ok, right?

Return values of type T are converted to sycl::ext::oneapi::experimental::uniform< T >, and broadcast to each work-item; every work-item in the sub-group receives the same value.

Copy link
Contributor Author

@sarnex sarnex Aug 8, 2023 •
edited
Loading

Choose a reason for hiding this comment

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

Yes, this is specifically needed because the invoke_simd spec defines uniform as a valid return type for ESIMD functions that are called through invoke_simd. We don't have a way to check at compile time if an ESIMD function is called by invoke_simd or not because it is/can be through a function pointer, so we just allow it in all cases.


void foo() {
constexpr unsigned Size = 1024;
constexpr unsigned GroupSize = 64;
sycl::range<1> GlobalRange{Size};
sycl::range<1> LocalRange{GroupSize};
sycl::nd_range<1> Range(GlobalRange, LocalRange);
queue q;
auto e = q.submit([&](handler &cgh) {
cgh.parallel_for(Range, [=](nd_item<1> ndi) {
uniform<int> x = invoke_simd(ndi.get_sub_group(), callee, 0);
});
});
}

int main() { foo(); }