-
Notifications
You must be signed in to change notification settings - Fork 787
[SYCL] Rename intel_reqd_sub_group_size attribute metadata #7713
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
This commit renames metadata nodes introduced by the intel_reqd_sub_group_size attribute feature design document from using intel_ prefix to use sycl_ prefix. Addresses intel#6824 Signed-off-by: Soumi Manna <[email protected]>
Follow-up comments from @steffenlarsen #6822 (comment) |
Signed-off-by: Soumi Manna <[email protected]>
Signed-off-by: Soumi Manna <[email protected]>
Signed-off-by: Soumi Manna <[email protected]>
I believe we will need the SPIR-V translator to also correctly consume the renamed metadata. @MrSidims - Can you confirm? |
@@ -9,7 +9,7 @@ kernel __attribute__((vec_type_hint(uint4))) __attribute__((work_group_size_hint | |||
// CHECK: define {{(dso_local )?}}spir_kernel void @kernel2(i32 {{[^%]*}}%a) {{[^{]+}} !vec_type_hint ![[MD3:[0-9]+]] !work_group_size_hint ![[MD4:[0-9]+]] | |||
|
|||
kernel __attribute__((intel_reqd_sub_group_size(8))) void kernel3(int a) {} | |||
// CHECK: define {{(dso_local )?}}spir_kernel void @kernel3(i32 {{[^%]*}}%a) {{[^{]+}} !intel_reqd_sub_group_size ![[MD5:[0-9]+]] | |||
// CHECK: define {{(dso_local )?}}spir_kernel void @kernel3(i32 {{[^%]*}}%a) {{[^{]+}} !sycl_reqd_sub_group_size ![[MD5:[0-9]+]] |
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.
I'm concerned about this change. This is an existing OpenCL C codegen test, why do we change the metadata name there? OpenCL flow is not being customized in our SYCL/DPC++ compiler as I understand it, so such change should go to upstream first, I suppose.
In OpenCL world, required sub-group size is an Intel extension and the translator is currently designed to only recognize intel_*
metadata. By changing its name we will effectively break support for the attribute. See corresponding test and related commit in the translator repo.
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.
I believe this change is a result of this discussion - #6804 (comment)
It isn't clear to me if the consensus was to actually modify this metadata. I don't see why we need to since this is an accepted Intel extension for openCL, and it exists outside of DPCPP. I understand the need to be consistent across various SYCL IR metadata but if SYCL is reusing existing constructs, I don't think we need to rename them. If we choose to do so, I agree with @AlexeySachkov that it should be done upstream.
I also can't recall why we have sycl::reqd_sub_group_size
, intel::reqd_sub_group_size
, intel_reqd_sub_group_size
and intel::sub_group_size
versions of this attribute. Can someone please explain?
IMO what is important right now is to identify all the |
Sorry for a long response, don't hesitate to ping me :) Yes, it would require the appropriate change. But I have the same concern as Alexey. |
This commit renames metadata nodes introduced by the intel_reqd_sub_group_size attribute feature design document from using intel_ prefix to use sycl_ prefix.
Addresses #6824
Signed-off-by: Soumi Manna [email protected]