[SYCL] Add SYCLPropagateAspectUsage Pass #5348
Conversation
llvm/include/llvm/SYCLPropagateAspectUsage/PropagateAspectUsage.h
Outdated
Show resolved
Hide resolved
llvm/test/SYCLLowerIR/PropagateAspectUsagePass/composite_types1.ll
Outdated
Show resolved
Hide resolved
llvm/test/SYCLLowerIR/PropagateAspectUsagePass/composite_types3.ll
Outdated
Show resolved
Hide resolved
llvm/test/SYCLLowerIR/PropagateAspectUsagePass/composite_types4.ll
Outdated
Show resolved
Hide resolved
|
Added consts where it is possible. In some places it is hard because it would require to add const inside some using aliases which will break the code. |
maksimsab
changed the title
| for (const Type *TT : T->subtypes()) { | ||
| // If TT = %A*** then we want to get TT = %A | ||
| // The same with arrays and vectors | ||
| while (TT->isPointerTy() || TT->isArrayTy() || TT->isVectorTy()) { |
There was a problem hiding this comment.
Typed pointers should be replaced by opaque pointers soon (https://discourse.llvm.org/t/opaque-pointers-status-update/60296/9, https://llvm.org/docs/OpaquePointers.html).
Could you please check that this code works for opaque pointers as well?
There was a problem hiding this comment.
I don't see that any of those APIs is going to be removed and therefore this code should continue to compile even with opaque pointers and it will continue to work correctly for arrays and vectors.
Pointer "branch" of this code won't work as it used to, because I assume there would be no contained type in it anymore.
@maksimsab, I think we should do the following:
| while (TT->isPointerTy() || TT->isArrayTy() || TT->isVectorTy()) { | |
| while ((TT->isPointerTy() && !TT->isOpaquePointerTy()) || TT->isArrayTy() || TT->isVectorTy()) { |
This will skip opaque pointers and they will be handled elsewhere when they are accessed and this code will work for non-opaque pointers as intended by looking through them.
There was a problem hiding this comment.
In fact, I realized that opaque pointers might break the current solution because the first step of algorithm (which traverses over type graph) becomes insufficient for overall correctness.
| #error "ASPECT macro is expected to be defined" | ||
| #endif | ||
|
|
||
| ASPECT(0, "host") |
There was a problem hiding this comment.
https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:device-aspects
Please, make sure this list is aligned with the spec.
I don't see 'host' and 'host_debuggable' aspects there, so I think these must start with "ext_". Right?
To make it easier to follow, please, keep the order of standard aspect the same as in the spec. atomic64 should be moved.
There was a problem hiding this comment.
My understanding is that this file should be aligned with our implementation in headers to match values assigned there - from that point of view, I would prefer to follow the same order as in headers simply not to miss anything.
Whether or not that order in headers should be changed, why it contains host and whether or not it is ok to re-order enum values (comparing to the spec) - those are questions to @intel/llvm-reviewers-runtime
llvm/sycl/include/CL/sycl/aspects.hpp
Lines 15 to 17 in c818709
There was a problem hiding this comment.
@maksimsab, but we still need to (at least temporarily) add a comment into sycl/include/CL/sycl/aspects.hpp saying that whenever this file is changed, the pass (Aspects.def) must be adjusted as well
| @@ -0,0 +1,50 @@ | |||
| // Test checks Propagate Aspect Usage pass. | |||
There was a problem hiding this comment.
Perhaps @AaronBallman can answer this as he reviewed https://reviews.llvm.org/D106030. I don't know what is the right directory for the tests checking diagnostics emitted from LLVM passes.
| // 6 is taken from sycl/include/CL/sycl/aspects.hpp | ||
| static constexpr int AspectFP64 = 6; |
There was a problem hiding this comment.
Do I understand it correctly that here we are injecting dependencies on external project to the compiler?
Imagine how hard it will be to debug bugs in this code if I update only sycl/include/CL/sycl/aspects.hpp.
There was a problem hiding this comment.
That is correct, both the compiler and the runtime must be in sync.
I don't really want to introduce a hard dependency between them and therefore I don't know what would be a good way to handle that. I suppose we should at least add a comment to aspects.hpp saying that if this file is changed, then PorpagateAspectUsage pass must also be changed, what do you think?
| Struct s; | ||
| s.a = 1; | ||
| return s.a; | ||
| } | ||
|
|
||
| [[sycl::device_has(aspect::fp64)]] int checkStructUsesAspect(int a) { |
There was a problem hiding this comment.
Unused function argument 'a' and for funcWithStruct
| } | ||
|
|
||
| // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o /dev/null 2>&1 | FileCheck %s --check-prefix CHECK-EMPTY | ||
| // CHECK-EMPTY-NOT: checkEmptyDiveceHas |
There was a problem hiding this comment.
I'm not sure what you are testing here. Can you please explain? Are you trying to test function is not generated?
There was a problem hiding this comment.
I was testing here that this function is not present in warning messages. I will fix it right now.
There was a problem hiding this comment.
Oh I see. You were testing warning was not being generated for empty argument list. I don't think you need to test that since that comes under semantic tests of device_has attribute which has already been tested in FE.
I think in this PR you just need to test propagation and usage of aspects. So for a function with device_has attribute with empty argument list, a good test will be to use a type with aspect in that function to ensure warning is generated
There was a problem hiding this comment.
FE changes LGTM for now. As discussed in PR, diagnostics should be fixed in a followup PR - #5877
This patch introduces a new IR pass from OptionalDeviceFeatures doc which propagates aspect usages from declared types to functions which somehow use these types.