[SYCL] Merge sycl_declared_aspects to sycl_used_aspects #7419
Conversation
To make RT be able to throw an exception if device doesn't support the feature pass should also propagate sycl_declared aspects by merging it with sycl_used_aspects. Spec: intel#7415
| @@ -8,25 +8,25 @@ queue q; | |||
|
|
|||
| // CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]] | |||
|
|
|||
| // CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] { | |||
| // CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] | |||
There was a problem hiding this comment.
Do we run the pass even if -disable-llvm-passes option is passed?
There was a problem hiding this comment.
I'm trying to understand why the clang-only test was changed :)
There was a problem hiding this comment.
Hm, I fixed this test because it was failed. Seems like we are.
There was a problem hiding this comment.
There was a problem hiding this comment.
Seems like we need this FE test changes for now
There was a problem hiding this comment.
I assume sycl_used_aspects metadata is being added here since the pass runs even if we pass -disable-llvm-passes. Can you modify the CHECK to test for it as well?
There was a problem hiding this comment.
Exactly. Sure thing
There was a problem hiding this comment.
Please modify the comment at the top of the test to indicate we check sycl_used_aspects here as well.
There was a problem hiding this comment.
@premanandrao Done 586281c
There was a problem hiding this comment.
I have a generic question about this change. sycl_declared_aspects metadata is used to denote what aspects a kernel supports and sycl_used_aspects is used to represent what the kernel actually uses right? If you merge the 2, aren't you losing information on what is actually used in the kernel? Am I misunderstanding something?
| @@ -8,25 +8,30 @@ queue q; | |||
|
|
|||
| // CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]] | |||
|
|
|||
| // CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] { | |||
| // CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] | |||
| // CHECK-SAME: !sycl_used_aspects | |||
There was a problem hiding this comment.
IIUC the aspect list should be the same as that in !sycl_decl_aspects right? Please add ![[ASPECT1]] to this CHECK as well, and make similar modifications to remaining CHECK lines.
There was a problem hiding this comment.
Added e647f36.
About your question. Since sycl_declared_aspects strictly requires the device to support declared aspects whether they are really used or no. I believe It's ok to merge them into sycl_used_aspects so we treat them as used. @steffenlarsen am I right?
There was a problem hiding this comment.
My point was that you lose information about which aspects were actually used in the kernel once this new pass runs. I am not sure for what all purposes actual aspect usage information is required, but off the top of my head I expect it is required to emit proper diagnostics, etc. I am assuming that this merge is done only after any pass which requires knowledge of actual aspects used is complete?
There was a problem hiding this comment.
Yes, we perform this merge at the end of the pass.
There was a problem hiding this comment.
Since the merge happens after the verification it should be okay, though you could argue that if functions disagree about aspects in their device_has attributes (or properties) it should also be a warning. That said, the spec only says a diagnostic needs to be issued if a kernel with a device_has "uses" an optional kernel feature.
After the pass, I don't see any problem in considering all device_has aspects as propagated uses as in some ways that is how the runtime will see them too. However, maybe there could be value in preserving information about the origin of the aspects, i.e. types/functions explicitly marked as using an aspect such as sycl::half. We could do this by changing the name of the metadata produced from the __sycl_detail__::__uses_aspects__ attribute, then in this pass we would extract that information and pass it along as we currently do the sycl_used_aspects. However, I don't know of a current use-case for this so I don't think it has to be part of these changes.
KornevNikita
requested review from
Fznamznon,
premanandrao and
elizabethandrews
|
Failures are unrelated: intel/llvm-test-suite#1380 |
Related compiler changes: intel/llvm#6989 and intel/llvm#7419
…llvm-test-suite#1314) Related compiler changes: intel#6989 and intel#7419
To make RT be able to throw an exception if device doesn't support the feature pass should also propagate sycl_declared aspects by merging it with sycl_used_aspects.
Design: #7415