Skip to content

[MatMul] Prolog build out, adding automatic swizzle generator for a few tile sizes #1900

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
Merged
merged 98 commits into from
Mar 20, 2023
Merged

[MatMul] Prolog build out, adding automatic swizzle generator for a few tile sizes #1900

merged 98 commits into from
Mar 20, 2023

Conversation

shmsong
Copy link

@shmsong shmsong commented Aug 10, 2022

This PR tries to build out matmul CTA tiling support, prioritizing Ampere ones at the moment.

Also in this PR is an inline bank conflict check utility that reports bank conflict on shared memory accesses.

Remaining TODO:

  • add asymmetric swizzle operators to support more irregular tile sizes.

shmsong and others added 30 commits July 11, 2022 22:15
@shmsong
use custom propagator in ampere TN
6f5ba21
@shmsong
add tile ordering utilities
2329caf
@shmsong
initial matmul scheduler implementation
121af43
@shmsong
use matmul scheduler prototype on ampere and turing test cases
f958c53
@shmsong
extend to support Volta
397f74c
@shmsong
minor cleanup
00d9a57
@shmsong
comment cleanup
d7035aa
@shmsong
minor fix
9ffc61d
@shmsong
add fragment iteration and use it in matmul scheduler
ed0f525
@shmsong
use scheduler params for tests
c972116
@shmsong
fragment support in double buffer
d12a90f
@shmsong
add register double buffering test cases
c306b9b
@shmsong
clean up custom transform propagator
63f561f
@shmsong
Merge remote-tracking branch 'origin/devel' into matmul_propagator
3d47c1f
@shmsong
rebase fix
29f88c7
@shmsong
comment
d029b9f
@shmsong
move bounded selector to common area
5ac053f
@shmsong
Add logic to handle fake boundary tensors in selection.
b51d247
@shmsong
naming and comment
aba5087
@shmsong
remove unused parameters from mma node
426c381
@shmsong
remove unnecessary parameters from mma ir node
6d4f377
@shmsong
rename scheduling variables
5e1f41f
@shmsong
change accumulator tv interface
1960da9
@shmsong @zasdfgbnm
Update torch/csrc/jit/codegen/cuda/scheduler/utils.h
3a411c2 
Co-authored-by: Gao, Xiang <[email protected]>
@shmsong
PR feedback
8f2e4da
@shmsong
Merge branch 'matmul_propagator' of https://github.com/csarofeen/pytorch
eef3a97 
 into matmul_propagator
@shmsong
pipe through parallel type position
6ad2967
@shmsong
Merge remote-tracking branch 'origin/devel' into matmul_propagator
65c8f0a
@shmsong
Revert "fragment support in double buffer"
cd03b00 
This reverts commit d12a90f.
@shmsong
Merge branch 'matmul_propagator' into fragment_iter
380dd66
@zasdfgbnm zasdfgbnm mentioned this pull request Oct 29, 2022
Merged
@zasdfgbnm zasdfgbnm mentioned this pull request Dec 18, 2022
Merged
@zasdfgbnm
Merge branch 'devel' of github.com:csarofeen/pytorch into matmul_swiz…
397f04e 
…zle_gen
@zasdfgbnm
Merge branch 'devel' of github.com:csarofeen/pytorch into matmul_swiz…
2587b96 
…zle_gen
@zasdfgbnm
Merge branch 'matmul_swizzle_gen' of github.com:csarofeen/pytorch int…
23a697c 
…o matmul_swizzle_gen
@zasdfgbnm
Merge branch 'devel' of github.com:csarofeen/pytorch into matmul_swiz…
039db32 
…zle_gen
@zasdfgbnm
fix
252ea22
@zasdfgbnm
Merge branch 'devel' of github.com:csarofeen/pytorch into matmul_swiz…
8b9f174 
…zle_gen
@zasdfgbnm
Merge branch 'matmul_swizzle_gen' of github.com:csarofeen/pytorch int…
c712456 
…o matmul_swizzle_gen
@zasdfgbnm
no gcd
271aa03
@zasdfgbnm
comment
cd93fcf
@zasdfgbnm
Merge branch 'devel' of github.com:csarofeen/pytorch into matmul_swiz…
8ff8bff 
…zle_gen
@zasdfgbnm
Merge branch 'devel' of github.com:csarofeen/pytorch into matmul_swiz…
d87b5ab 
…zle_gen
@zasdfgbnm
cleanup
f092dd6
@zasdfgbnm
Merge branch 'devel' of github.com:csarofeen/pytorch into matmul_swiz…
151d4be 
…zle_gen
@zasdfgbnm
Merge branch 'devel' of github.com:csarofeen/pytorch into matmul_swiz…
d04cd5b 
…zle_gen
@zasdfgbnm zasdfgbnm mentioned this pull request Mar 11, 2023
Merged
1 task
@zasdfgbnm zasdfgbnm mentioned this pull request Mar 20, 2023
Merged
1 task
@zasdfgbnm zasdfgbnm changed the base branch from devel to rebase-matmul_swizzle_gen March 20, 2023 17:56
@zasdfgbnm zasdfgbnm merged commit c8c8cd7 into rebase-matmul_swizzle_gen Mar 20, 2023
@zasdfgbnm zasdfgbnm deleted the matmul_swizzle_gen branch March 20, 2023 17:57
zasdfgbnm added a commit to NVIDIA/Fuser that referenced this pull request Mar 21, 2023
@zasdfgbnm
Fix dtype calculation in expr simplifier (#36)
623c049 
When working on csarofeen/pytorch#1900, I find
that sometimes expr simplifier will assign the dtype of a value wrongly.
This is because expr simplifier is written loosely assuming the dtypes
of all `Val`s are the same. However, because we are putting pointer
types into the expression, we need to be more careful, otherwise we will
get wrong kernel code like:
```
__half* ptr1 = threadIdx.x * 128;
__half* ptr2 = ptr1 + 256 + T1.data;
```
This PR changes all passes in the expr simplifier to let it infer dtype
from its inputs.

- [x] TODO: update all the failing `assertCUDAKernel` tests😵‍💫😵‍💫😵‍💫😵‍💫
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@naoyam naoyam naoyam left review comments

@zasdfgbnm zasdfgbnm zasdfgbnm left review comments

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
3 participants
@shmsong @naoyam @zasdfgbnm