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.

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

Conversation

sarnex
Copy link
Contributor

@sarnex sarnex commented Aug 7, 2023

The spec states that returning a uniform object is allowed:

"Return values of type sycl::ext::oneapi::experimental::uniform are not anyhow converted, and broadcast to each work-item; every work-item in the sub-group receives the same value. NOTE: sycl::ext::oneapi::experimental::uniform return type is the way to return a uniform value of simd or simd_mask type."

Update the compile-time error checking and ESIMD verifier to allow this.

The spec states that returning a `uniform` object is allowed:

"Return values of type sycl::ext::oneapi::experimental::uniform<T> are not anyhow converted, and broadcast to each work-item;
every work-item in the sub-group receives the same value.
NOTE: sycl::ext::oneapi::experimental::uniform<T> return type is the way to return a uniform value of simd or simd_mask type."

Update the compile-time error checking and ESIMD verifier to allow this.

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex sarnex temporarily deployed to aws August 7, 2023 15:25 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws August 7, 2023 16:29 — with GitHub Actions Inactive
@sarnex sarnex marked this pull request as ready for review August 7, 2023 17:48
@sarnex sarnex requested review from a team and rolandschulz as code owners August 7, 2023 17:48
Copy link
Contributor

@turinevgeny turinevgeny left a comment

Choose a reason for hiding this comment

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

LGTM.

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

@v-klochkov v-klochkov Aug 7, 2023

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

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.

@v-klochkov v-klochkov merged commit 5eef8c7 into intel:sycl Aug 8, 2023
mdtoguchi pushed a commit to mdtoguchi/llvm that referenced this pull request Oct 18, 2023
The spec states that returning a `uniform` object is allowed:

"Return values of type sycl::ext::oneapi::experimental::uniform<T> are
not anyhow converted, and broadcast to each work-item; every work-item
in the sub-group receives the same value. NOTE:
sycl::ext::oneapi::experimental::uniform<T> return type is the way to
return a uniform value of simd or simd_mask type."

Update the compile-time error checking and ESIMD verifier to allow this.

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants