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.

Sign up for GitHub

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

Jump to bottom

SdpaFwdOp::toString doesn't print all outputs. #3037

Closed
wujingyue opened this issue Sep 26, 2024 · 1 comment · Fixed by #3328
Closed

SdpaFwdOp::toString doesn't print all outputs. #3037

wujingyue opened this issue Sep 26, 2024 · 1 comment · Fixed by #3328
Assignees
@Priya2698
Labels
Matmuls Top-Down Matmul Dev

Comments

@wujingyue
Copy link
Collaborator

Fuser/csrc/ir/nodes.cpp

Line 4348 in 29a59a0

indent(ss, indent_size) << attn_out()->toString() << "\n";

This makes debugging harder because I was seeing a fusion output TensorView defined by nothing.
@wujingyue
Copy link
Collaborator Author

For example,

T55_g___bfloat[ bS163{1}, iS164{96}, iS165{2048}, iS166{128} ]
   = sdpa(T52_l___bfloat[ bS149{1}, iS151{96}, iS150{2048}, iS152{( ceilDiv(12288, 96) )} ],
            T50_l___bfloat[ bS139{1}, iS141{96}, iS140{2048}, iS142{( ceilDiv(12288, 96) )} ],
            T54_l___bfloat[ bS159{1}, iS161{96}, iS160{2048}, iS162{( ceilDiv(12288, 96) )} ],
            dropout_p = double(0.10000000000000001),
            is_causal = true  )

T56 is also defined by this Expr but not being printed.

@Priya2698 Priya2698 mentioned this issue Nov 1, 2024
Merged
@Priya2698 Priya2698 linked a pull request Nov 1, 2024 that will close this issue
Print all outputs in SdpaFwdOp #3328
Merged
wujingyue pushed a commit that referenced this issue Nov 2, 2024
@Priya2698
Print all outputs in SdpaFwdOp (#3328)
e6b285b 
Closes Issue #3037.

For the test:
https://github.com/NVIDIA/Fuser/blob/f08bd5182df894019cfd07f04f7b2125750af713/tests/cpp/test_sdpa_node.cpp#L691


**`SdpaFwdOp::toString()`**
```
T3_l___half[ iS12{16}, iS13{32}, iS14{64}, iS15{64} ],
T4_l_float[ iS16{16}, iS17{32}, iS18{64} ],
T5_l_int64_t[ ],
T6_l_int64_t[ ]
   = sdpa(T0_g___half[ iS0{16}, iS1{32}, iS2{64}, iS3{64} ],
            T1_g___half[ iS4{16}, iS5{32}, iS6{128}, iS7{64} ],
            T2_g___half[ iS8{16}, iS9{32}, iS10{128}, iS11{64} ],
            dropout_p = double(0),
            is_causal = false  )
```
**`SdpaBwdOp::toString()`**
```
T8_g___half[ iS23{16}, iS24{32}, iS25{64}, iS26{64} ],
T9_g___half[ iS27{16}, iS28{32}, iS29{128}, iS30{64} ],
T10_g___half[ iS31{16}, iS32{32}, iS33{128}, iS34{64} ]
   = sdpa_bwd(T7_g___half[ iS19{16}, iS20{32}, iS21{64}, iS22{64} ],
            T0_g___half[ iS0{16}, iS1{32}, iS2{64}, iS3{64} ],
            T1_g___half[ iS4{16}, iS5{32}, iS6{128}, iS7{64} ],
            T2_g___half[ iS8{16}, iS9{32}, iS10{128}, iS11{64} ],
            T3_g___half[ iS12{16}, iS13{32}, iS14{64}, iS15{64} ],
            logsum_exp = T4_l_float[ iS16{16}, iS17{32}, iS18{64} ],
            dropout_p = double(0),
            is_causal = false,
            philox_seed = T5_l_int64_t[ ],
            philox_offset = T6_l_int64_t[ ],
  )
```
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees

@Priya2698 Priya2698

Labels
Matmuls Top-Down Matmul Dev
Projects
None yet
Milestone
No milestone
Development

Successfully merging a pull request may close this issue.

Print all outputs in SdpaFwdOp NVIDIA/Fuser
3 participants
@wujingyue @kevinstephano @Priya2698