-
Notifications
You must be signed in to change notification settings - Fork 787
[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
Conversation
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>
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM.
[[intel::device_indirectly_callable]] uniform<int> | ||
callee(simd<int, 8>) SYCL_ESIMD_FUNCTION { | ||
return uniform<int>(5); | ||
} |
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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.
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>
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.