-
Notifications
You must be signed in to change notification settings - Fork 246
Fix TopologicalSort for Joint and Cooperative matrices #2143
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
Conversation
Signed-off-by: Sidorov, Dmitry <[email protected]>
This reverts commit 0330406.
visit method of the sort relies on getNonLiteralOperands method of the SPIRVType which is being inserted in the module. Without it dependent types can be inserted in the module in incorrect order. Signed-off-by: Sidorov, Dmitry <[email protected]>
| @@ -0,0 +1,436 @@ | |||
| ;; Compiled from joint_matrix_bf16_fill_k_cache.cpp from https://github.com/intel/llvm | |||
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.
Honestly this test can be much smaller, like:
1 ; ModuleID = 'artificial.bc'
2 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
3 target triple = "spir64-unknown-unknown"
4
5 %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix" = type { target("spirv.CooperativeMatrixKHR", float, 8, 16, 0, 0) }
6
7 $matrix_fun = comdat any
8
9 ; Function Attrs: convergent norecurse nounwind
10 define weak_odr dso_local spir_kernel void @matrix_fun() {
11 entry:
12 %tC.i = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix"
13 ret void
14 }
but since I've removed many matrix test during switch to opaque pointers with an idea to re-generate them later with the appropriate FE, I've decided that may be this patch is a good casus to start doing it :)
and since joint/coop matrix is quite complicated feature it's nice to have real-life IR modules being translated
visit method of the sort relies on getNonLiteralOperands method of the SPIRVType which is being inserted in the module. Without it dependent types can be inserted in the module in incorrect order. For example: TypeCooperativeMatrixKHR %ID% TypeStruct ... %ID% is the correct order, but without the patch in some cases the translator could generate the opposite order. It's a cherry-pick of KhronosGroup/SPIRV-LLVM-Translator#2143 Signed-off-by: Sidorov, Dmitry <[email protected]>
…2143) visit method of the sort relies on getNonLiteralOperands method of the SPIRVType which is being inserted in the module. Without it dependent types can be inserted in the module in incorrect order. For example: TypeCooperativeMatrixKHR %ID% TypeStruct ... %ID% is the correct order, but without the patch in some cases the translator could generate the opposite order. Signed-off-by: Sidorov, Dmitry <[email protected]> (cherry picked from commit 436c497)
…2143) visit method of the sort relies on getNonLiteralOperands method of the SPIRVType which is being inserted in the module. Without it dependent types can be inserted in the module in incorrect order. For example: TypeCooperativeMatrixKHR %ID% TypeStruct ... %ID% is the correct order, but without the patch in some cases the translator could generate the opposite order. Signed-off-by: Sidorov, Dmitry <[email protected]> (cherry picked from commit 436c497)
…2143) visit method of the sort relies on getNonLiteralOperands method of the SPIRVType which is being inserted in the module. Without it dependent types can be inserted in the module in incorrect order. For example: TypeCooperativeMatrixKHR %ID% TypeStruct ... %ID% is the correct order, but without the patch in some cases the translator could generate the opposite order. Signed-off-by: Sidorov, Dmitry <[email protected]> (cherry picked from commit 436c497)
…2143) visit method of the sort relies on getNonLiteralOperands method of the SPIRVType which is being inserted in the module. Without it dependent types can be inserted in the module in incorrect order. For example: TypeCooperativeMatrixKHR %ID% TypeStruct ... %ID% is the correct order, but without the patch in some cases the translator could generate the opposite order. Signed-off-by: Sidorov, Dmitry <[email protected]> (cherry picked from commit 436c497)
…atrices (KhronosGroup#2143) visit method of the sort relies on getNonLiteralOperands method of the SPIRVType which is being inserted in the module. Without it dependent types can be inserted in the module in incorrect order. For example: TypeCooperativeMatrixKHR %ID% TypeStruct ... %ID% is the correct order, but without the patch in some cases the translator could generate the opposite order. Signed-off-by: Sidorov, Dmitry <[email protected]> (cherry picked from commit 436c497)
…atrices (#2143) (#3242) visit method of the sort relies on getNonLiteralOperands method of the SPIRVType which is being inserted in the module. Without it dependent types can be inserted in the module in incorrect order. For example: TypeCooperativeMatrixKHR %ID% TypeStruct ... %ID% is the correct order, but without the patch in some cases the translator could generate the opposite order. Signed-off-by: Sidorov, Dmitry <[email protected]> (cherry picked from commit 436c497) Signed-off-by: Sidorov, Dmitry <[email protected]> Co-authored-by: Dmitry Sidorov <[email protected]>
visit method of the sort relies on getNonLiteralOperands method of the
SPIRVType which is being inserted in the module. Without it dependent
types can be inserted in the module in incorrect order.
For example:
TypeCooperativeMatrixKHR %ID%
TypeStruct ... %ID%
is the correct order, but without the patch in some cases the translator
could generate the opposite order.
Signed-off-by: Sidorov, Dmitry [email protected]