[fp8 rowwise] Support non-contiguous inputs and clarify checks #134225
Conversation
🔗 Helpful Links🧪 See artifacts and rendered test results at hud.pytorch.org/pr/134225
Note: Links to docs will display an error until the docs builds have been completed. ❌ 1 New FailureAs of commit ffec533 with merge base aea1148 ( NEW FAILURE - The following job has failed:
This comment was automatically generated by Dr. CI and updates every 15 minutes. |
|
|
||
| StrideInputA stride_a = cutlass::make_cute_packed_stride( | ||
| StrideInputA{}, cute::make_shape(M, K, 1)); | ||
| StrideInputA{}, cute::make_shape(M, static_cast<int>(XQ.stride(0)), 1)); |
There was a problem hiding this comment.
Also total noob Q / more curiosity the trailing 1 is referring to this right: https://github.com/NVIDIA/cutlass/blob/f7b19de32c5d1f3cedfc735c2849f12b537522ee/include/cutlass/gemm/kernel/sm90_gemm_warpspecialized_cooperative.hpp#L321-L323?
There was a problem hiding this comment.
Uhm, no, it's just saying that the innermost dimension (i.e., K) is contiguous. It would be equivalent to put XQ.stride(1) here, but since we have previously checked that this equals 1, I just hardcoded 1 right away.
There was a problem hiding this comment.
This is cutlass noob Q I guess I am suprrised that the strides appear to be of len 3 instead of 2, e.g. M, 1, versus (M, 1, 1)
It seems the trailing dim is refers to a L (local tile) ?
It's likely a matter of opinion, but I find this new version to have less duplication, even if it might have more boilerplate. Pull Request resolved: #134226 Approved by: https://github.com/drisspg ghstack dependencies: #134223, #134224, #134225
Pull Request resolved: #134227 Approved by: https://github.com/drisspg, https://github.com/eqy ghstack dependencies: #134223, #134224, #134225, #134226
It's likely a matter of opinion, but I find this new version to have less duplication, even if it might have more boilerplate. Pull Request resolved: pytorch#134226 Approved by: https://github.com/drisspg ghstack dependencies: pytorch#134223, pytorch#134224, pytorch#134225
) Pull Request resolved: pytorch#134227 Approved by: https://github.com/drisspg, https://github.com/eqy ghstack dependencies: pytorch#134223, pytorch#134224, pytorch#134225, pytorch#134226
3 participants
Stack from ghstack (oldest at bottom):