Skip to content
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

Merged
merged 8 commits into from
Jan 20, 2025

Conversation

MrSidims
Copy link
Contributor

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.

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]>
@MrSidims MrSidims requested review from a team as code owners January 14, 2025 14:58
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) {
Copy link
Contributor

@dkhaldi dkhaldi Jan 14, 2025

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

Copy link
Contributor Author

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]>
@MrSidims
Copy link
Contributor Author

Thanks for the comments, I have applied them.

Copy link
Contributor

@dkhaldi dkhaldi left a 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]>
@MrSidims MrSidims requested a review from YuriPlyakhin January 16, 2025 16:23
Copy link
Contributor

@YuriPlyakhin YuriPlyakhin left a 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.

@MrSidims MrSidims requested a review from a team January 17, 2025 11:28
@MrSidims
Copy link
Contributor Author

@intel/dpcpp-tools-reviewers @AlexeySachkov friendly ping

@asudarsa
Copy link
Contributor

@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(
Copy link
Contributor

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?

Copy link
Contributor Author

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

Copy link
Contributor

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.

Copy link
Contributor

@asudarsa asudarsa left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@MrSidims
Copy link
Contributor Author

@intel/llvm-gatekeepers please help with the merge

@MrSidims MrSidims requested a review from a team January 20, 2025 12:00
@dm-vodopyanov dm-vodopyanov merged commit e257292 into intel:sycl Jan 20, 2025
17 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

6 participants