-
Notifications
You must be signed in to change notification settings - Fork 755
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
[SYCL][Matrix] Propagate constexpr matrix layout even with O0 #16628
Conversation
Per SPIR-V specification Layout of a matrix must be a constant instruction aka a constexpr or specialization constant. Meanwhile in SYCL headers layout is passed as a parameter to joint_matrix_load function, so even if that layout is a constant expression in the user's code - it's not possible to prove that to the compiler, so constant propagation will happen only after inlining, not in AST. That means, that with O0 layout would remain to be a runtime variable in LLVM IR. SYCL matrix layout is being mapped on SPIR-V matrix layout by joint_matrix_layout_to_spv function. This patch adds routine that finds calls to this function and replaces them with the found constant. To help this routine always_inline attribute was removed from joint_matrix_layout_to_spv function. Signed-off-by: Sidorov, Dmitry <[email protected]>
Signed-off-by: Sidorov, Dmitry <[email protected]>
@@ -68,7 +68,7 @@ convertMatrixUseStringToEnum(const char *UseString) { | |||
return std::nullopt; | |||
} | |||
|
|||
inline __SYCL_ALWAYS_INLINE __spv::MatrixLayout joint_matrix_layout_to_spv( | |||
constexpr inline __spv::MatrixLayout joint_matrix_layout_to_spv( | |||
sycl::ext::oneapi::experimental::matrix::layout Layout) { |
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 guess it was not possible to capture and preserve the layout as constant at this stage?
That's why you had to create this preservation as a function in SYCLJointMatrixTransform
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.
Compiler considers it to be a runtime value as joint_matrix_load/store
functions (that calls joint_matrix_layout_to_spv
) can not be in constexpr
context, as they call external functions (__spirv_CooperativeMatrixLoad/StoreKHR
). I'm leaving constexpr
keyword here as it's really our intention, yet not enforced by the compiler.
Signed-off-by: Sidorov, Dmitry <[email protected]>
Signed-off-by: Sidorov, Dmitry <[email protected]>
Thanks for the comments, I have applied them. |
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
Signed-off-by: Sidorov, Dmitry <[email protected]>
Signed-off-by: Sidorov, Dmitry <[email protected]>
Signed-off-by: Sidorov, Dmitry <[email protected]>
Signed-off-by: Sidorov, Dmitry <[email protected]>
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. One thought of alternative algorithm in the comments, but no strong opinion.
@intel/dpcpp-tools-reviewers @AlexeySachkov friendly ping |
Answered above. Thanks |
inline __SYCL_ALWAYS_INLINE __spv::MatrixLayout joint_matrix_layout_to_spv( | ||
// propagateConstexprLayout uses the exact name of the function, so we use | ||
// extern "C" here. | ||
extern "C" constexpr __spv::MatrixLayout joint_matrix_layout_to_spv( |
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.
Do we not need the inlining attribute anymore?
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.
No, as constexpr
also includes inline
, see the comment from Alexey
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.
Ah...good...sorry..missed that as it was hiding.
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/llvm-gatekeepers please help with the merge |
Per SPIR-V specification Layout of a matrix must be a constant instruction aka a constexpr or specialization constant. Meanwhile in SYCL headers layout is passed as a parameter to joint_matrix_load function, so even if that layout is a constant expression in the user's code - it's not possible to prove that to the compiler, so constant propagation will happen only after inlining, not in AST. That means, that with O0 layout would remain to be a runtime variable in LLVM IR.
SYCL matrix layout is being mapped on SPIR-V matrix layout by joint_matrix_layout_to_spv function. This patch adds routine that finds calls to this function and replaces them with the found constant. To help this routine always_inline attribute was removed from joint_matrix_layout_to_spv function.