[SYCL][clang] Fix uses_aspects applied to function declarations #10164
Conversation
It was only applied to definitions.
|
@elizabethandrews , could you please review? |
| @@ -38,6 +38,9 @@ constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; } | |||
| // CHECK: define dso_local spir_func void @{{.*}}func7{{.*}} !sycl_used_aspects ![[ASPECTS1]] { | |||
| [[__sycl_detail__::__uses_aspects__(getAspect())]] void func7() {} | |||
|
|
|||
| // CHECK: declare !sycl_used_aspects ![[ASPECTS1]] spir_func void @{{.*}}func8{{.*}} | |||
| [[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] SYCL_EXTERNAL void func8(); | |||
There was a problem hiding this comment.
The design document specifies this metadata is applied to function definitions - https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/doc/design/OptionalDeviceFeatures.md
How does this work with redeclarations? What is finally attached to the function definition? IIRC there is some pass which propagates this metadata up the static call graph. What gets propagated now if we have different attribute arguments in different declarations
There was a problem hiding this comment.
The design document specifies this metadata is applied to function definitions - https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/doc/design/OptionalDeviceFeatures.md
The document seems to be assuming that the atribute can be applied to function declarations, see for example https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/doc/design/OptionalDeviceFeatures.md#changes-to-dpc-headers . However it says,
If a function is decorated with [[sycl_detail::uses_aspects()]], the front-end adds an !sycl_used_aspects metadata to the function's definition listing the aspects from that attribute.
If a function is decorated with [[sycl::device_has()]], the front-end adds an !sycl_declared_aspects metadata to the function's definition listing the aspects from that attribute.
For both uses_aspects and device_has attibutes. We already has support for device_has on declarations added by https://github.com/intel/llvm/pull/9611/files . I'm not sure why it is limited to definitions in the document, maybe we should change the wording to "function signature"? CC @gmlueck, @AlexeySachkov to clarify.
How does this work with redeclarations? IIRC there is some pass which propagates this metadata up the static call graph. What gets propagated now if we have different attribute arguments in different declarations
When redeclarations appear, the most recent is referenced by the code. Meaning the metadata is generated based on attirbute value applied to the most recent declarations. Even in case like
SYCL_EXTERNAL [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp64)]] void foo() ;
kernel1 calling foo
SYCL_EXTERNAL [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp16)]] void foo() ;
kernel2 calling foo
I'm seeing only fp16 aspect on kernels and function declaration.
The pass seem to be updated by https://github.com/intel/llvm/pull/9611/files in generic way, so it seems to be working fine with declarations.
The metadata generated for the most recent redeclaration is propagated. When the most recent doesn't have the attribute, metadata is propagated from previous redeclaration.
The thing I find confusing for the end user, when attribute is applied to a redeclaration and the previous declaration had an attribute, the warning is emitted:
t.cpp:29:17: warning: attribute '__uses_aspects__' is already applied [-Wignored-attributes]
29 | SYCL_EXTERNAL [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp16)]] void foo();
| ^
t.cpp:4:17: note: previous attribute is here
4 | SYCL_EXTERNAL [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp64)]] void foo()
| ^
1 warning generated.
It seems to be saying that the attribute is ignored on the most recent declaration, however instead it seems to be ignored on the first declaration since the metadata is generated based on the latter declaration.
There was a problem hiding this comment.
The [[sycl::device_has]] attribute is defined in the SYCL 2020 specification section 5.8 "Attributes for device code", which says this about attributes and redeclarations:
If one of the attributes defined in this section is applied to a kernel or device function, it must be applied
to the first declaration of that kernel or device function in the translation unit. Programs which fail to do
this are ill formed and the compiler must issue a diagnostic. Redeclarations of the kernel or device function in the same translation unit may optionally have the same attribute applied (so long as the attribute
arguments are the same between the declarations), but this is not required. The attribute remains in
effect regardless of whether it appears in the redeclaration.
Based on that wording, I think it is an error to redeclare a function with a different [[sycl::device_has]] attribute, and the compiler is supposed to diagnose this case.
I think it is reasonable to treat the [sycl_detail::uses_aspects]] attribute in the same way.
BTW, the links you have above are pointing to a different company's fork of our repo. When looking at our design documents, you should look at our own fork (not the triSYCL fork). For example, this is the link to our copy of the optional device features design document:
https://github.com/intel/llvm/blob/sycl/sycl/doc/design/OptionalDeviceFeatures.md
There was a problem hiding this comment.
BTW, the links you have above are pointing to a different company's fork of our repo. When looking at our design documents, you should look at our own fork (not the triSYCL fork)
OMG, I didn't notice that the first link posted by Elizabeth is to triSYCL repo. Thanks for pointing this out.
Based on that wording, I think it is an error to redeclare a function with a different [[sycl::device_has]] attribute, and the compiler is supposed to diagnose this case.
Well, apparently the compiler's behavior is different now. I would prefer implementing the diagnostic separately though.
There was a problem hiding this comment.
BTW, the links you have above are pointing to a different company's fork of our repo.
I didn't notice I checked the wrong link! I just opened the first search item on google. I apologize!
There was a problem hiding this comment.
It seems to be saying that the attribute is ignored on the most recent declaration, however instead it seems to be ignored on the first declaration since the metadata is generated based on the latter declaration.
I believe this is 'normal' clang attribute functionality. Warnings are usually generated on redeclarations but metadata is also usually from first declaration. So there is some bug there. Anyway based on the spec @gmlueck pasted above, it looks like we need to be parsing attribute arguments and generating this diagnostic accordingly. I agree that can be a separate patch since it is orthogonal to this PR.
|
Thank you @elizabethandrews . I'll create a tracker to correct diagnostic according to the spec. |
It was only applied to definitions.