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

Enable transpose scheduler #1927

Merged

zasdfgbnm merged 48 commits into devel from enable-transpose-scheduler

Sep 11, 2022

Collaborator

zasdfgbnm commented Aug 24, 2022 •

edited

Loading

transpose benchmarks (RTX 3090)

Command:

./build/bin/nvfuser_bench --benchmark_out=filename.csv --benchmark_out_format=csv "--benchmark_filter=.*Transpose.*

Raw data (before): https://github.com/zasdfgbnm/things/blob/master/2022/nvfuser-benchmarks/transpose-pointwise.csv
Raw data (after): https://github.com/zasdfgbnm/things/blob/master/2022/nvfuser-benchmarks/transpose-transpose.csv

Geomean	1.14
max	12.08
top 1%	9.26
top 10%	1.43
top 50%	1.00
top 90%	0.97
top 99%	0.96
min	0.95

Benchmarks of the entire suite (RTX 3090)

Geomean	1.02
max	12.08
top 1%	2.51
top 10%	1.01
top 50%	1.00
top 90%	0.98
top 99%	0.96
min	0.72

zasdfgbnm and others added 10 commits

August 24, 2022 10:19


Enable transpose scheduler

e84fff8


add lower index resolution

bb8cef1


Merge remote-tracking branch 'origin/devel' into local_index_expr

66765a7


add repro


clear GPU memory after test

e9d09fe


cleanup and comment

f406e23

fix

abe1f6d

fix

bcc6f6c


Allow splitting inner-most ID to create virtual innermost ID

63e3e76


remove obselete comment

5a423a0

zasdfgbnm changed the base branch from devel to local_index_expr

August 26, 2022 03:37

zasdfgbnm changed the base branch from local_index_expr to devel

August 26, 2022 03:40

zasdfgbnm added 3 commits

August 25, 2022 20:40


Merge branch 'local_index_expr' of github.com:csarofeen/pytorch into …

0b30aa1

…enable-transpose-scheduler


Merge branch 'transpose-split-inner' into enable-transpose-scheduler

8bfaa28


Merge branch 'devel' of github.com:csarofeen/pytorch into enable-tran…

217eb96

…spose-scheduler

zasdfgbnm changed the base branch from devel to transpose-split-inner

August 27, 2022 07:15

zasdfgbnm added 14 commits

August 27, 2022 00:15


Merge branch 'devel' of github.com:csarofeen/pytorch into transpose-s…

bf956ac

…plit-inner


Merge branch 'transpose-split-inner' into enable-transpose-scheduler

b796975


Merge branch 'enable-transpose-scheduler' of github.com:csarofeen/pyt…

10f64f5

…orch into enable-transpose-scheduler


Merge branch 'devel' of github.com:csarofeen/pytorch into enable-tran…

288b131

…spose-scheduler


Merge branch 'devel' of github.com:csarofeen/pytorch into transpose-s…

…plit-inner


Merge branch 'transpose-split-inner' into enable-transpose-scheduler

3752e80


Merge branch 'devel' of github.com:csarofeen/pytorch into transpose-s…

…plit-inner


Merge branch 'transpose-split-inner' into enable-transpose-scheduler

090b582


Merge branch 'devel' of github.com:csarofeen/pytorch into transpose-s…

27b6893

…plit-inner


Merge branch 'transpose-split-inner' into enable-transpose-scheduler

f506e87


skip innermost split

5f266b3


comment

bd93302


cleanup

5933f53


save

7c366b7

Base automatically changed from transpose-split-inner to devel

September 6, 2022 19:51

zasdfgbnm added 2 commits

September 6, 2022 13:12


Merge branch 'devel' of github.com:csarofeen/pytorch into enable-tran…

ab0c880

…spose-scheduler


reject < 1 wave on runtime

21a7896

jjsjann123 mentioned this pull request

[nvfuser] view size is not compatible with input tensor's size and stride pytorch/pytorch#93648

Closed

zasdfgbnm added 2 commits

September 8, 2022 15:19


cache inner most positions

f8551d7


issue id

53690a0

zasdfgbnm commented

View reviewed changes

torch/csrc/jit/codegen/cuda/scheduler/transpose.cpp

Comment on lines +228 to +231

+ if (wave_elements >= n_elems) {
+ // if one full wave can handle all elements, don't create virtual inner dims
+ return;
+ }

Collaborator Author

zasdfgbnm Sep 8, 2022

This case should be rejected in canScheduleRunTime, but I think it is not a bad idea to still handle it inside the heuristics.

zasdfgbnm commented

View reviewed changes

torch/csrc/jit/codegen/cuda/scheduler/transpose.cpp

Comment on lines -339 to -349

-bool hasAtLeastTwoValidGroups(Fusion* fusion) {
- return DomainMap::hasAtLeastTwoValidGroups(fusion);
-}
-std::shared_ptr<TransposeParams> getTransposeHeuristics(
- Fusion* fusion,
- const at::ArrayRef<c10::IValue>& runtime_inputs,
- HeuristicSummary* data_cache) {
- SchedulerRuntimeInfo runtime_info(fusion, runtime_inputs, true);
- return getTransposeHeuristics(fusion, runtime_info, data_cache);
-}

Collaborator Author

zasdfgbnm Sep 8, 2022

No real change, just moved below

zasdfgbnm commented

View reviewed changes

torch/csrc/jit/codegen/cuda/scheduler/transpose.cpp Outdated

- // Incase any buffer is of type DataType::Index
- DataType index_type = indexModeToDtype(runtime_info.getIndexMode());
+HeuristicSummaryEntry<HeuristicCompileTime::DomainMap> getDomainMap(

Collaborator Author

zasdfgbnm Sep 8, 2022

Split-out from getTransposeHeuristics for code reuse. Should have no behavioral change.

zasdfgbnm commented

View reviewed changes

torch/csrc/jit/codegen/cuda/scheduler/transpose.cpp


		HeuristicSummaryEntry<HeuristicCompileTime::InputsOutputsInnerDimGroups>
		getInputsOutputsGroups(HeuristicSummary* data_cache, DomainMap& domain_map) {

Collaborator Author

zasdfgbnm Sep 8, 2022

Split-out from getTransposeHeuristics for code reuse. Should have no behavioral change.

zasdfgbnm commented

View reviewed changes

torch/csrc/jit/codegen/cuda/scheduler/transpose.cpp

+}
+HeuristicSummaryEntry<HeuristicCompileTime::ReferenceTensors>
+getReferenceTensors(

Collaborator Author

zasdfgbnm Sep 8, 2022

Split-out from getTransposeHeuristics for code reuse. Should have no behavioral change.

zasdfgbnm commented

View reviewed changes

torch/csrc/jit/codegen/cuda/scheduler/transpose.cpp

- auto ref_root = reference1->getMaybeRFactorDomain();
- std::vector<int64_t> shape_in_ref1;
- shape_in_ref1.reserve(reference1->nDims());
+std::pair<std::vector<int64_t>, int64_t> getShapeInReference(

Collaborator Author

zasdfgbnm Sep 8, 2022

Split-out from getTransposeHeuristics for code reuse. Should have no behavioral change.

zasdfgbnm commented

View reviewed changes

torch/csrc/jit/codegen/cuda/scheduler/transpose.cpp

- auto params = std::make_shared<TransposeParams>("Transpose heuristics");
+HeuristicSummaryEntry<HeuristicCompileTime::InnerMostDimInfo>
+getInnerMostDimInfoInReference(

Collaborator Author

zasdfgbnm Sep 8, 2022

The position of innermost dim is now cached.


lint

44f7810

zasdfgbnm changed the title ~~[WIP] Enable transpose scheduler~~ Enable transpose scheduler

Collaborator Author

zasdfgbnm commented Sep 8, 2022

@csarofeen I am marking this as ready, but I still need to rerun the benchmarks after fixing the effective bandwidth calculation.

zasdfgbnm marked this pull request as ready for review

September 8, 2022 23:54

zasdfgbnm added 6 commits

September 8, 2022 21:27


Merge branch 'devel' of github.com:csarofeen/pytorch into enable-tran…

98ceddd

…spose-scheduler


maybe_tune

ea0d1ff


Revert "maybe_tune"

5920c62

This reverts commit ea0d1ff.


reject at pointwise scheduler

fix

27a4db8

fix

0c173fa

zasdfgbnm commented

View reviewed changes

torch/csrc/jit/codegen/cuda/scheduler/registry.cpp

@@ @@ -937,6 +937,84 @@ class ReductionScheduler : public SchedulerEntry { @@
  }
 };
+class TransposeScheduler : public SchedulerEntry {

Collaborator Author

zasdfgbnm Sep 11, 2022

moved from below


fix fusion input reductions

911956e

Collaborator Author

zasdfgbnm commented Sep 11, 2022

Should be good to go. Please check updated benchmarks #1927 (comment)

csarofeen approved these changes

View reviewed changes

Owner

csarofeen left a comment

LGTM

zasdfgbnm merged commit bd93578 into devel

zasdfgbnm deleted the enable-transpose-scheduler branch

September 11, 2022 22:09

jjsjann123 added a commit that referenced this pull request


[NVFuser] Upstream push 1026 (pytorch#87779)

7b419e8

Syncing nvfuser devel branch to upstream master. https://github.com/csarofeen/pytorch/

Codegen changes include:

* codegen improvement:
    i. allow non-root trivial reductions, allow empty/no-op fusion
    ii. fixes vectorization checks and size calculation
    iii. bank conflict handle improvement
    iv. enables transpose scheduler

* misc:
    i. CI tests failure fixes
    ii. cpp tests file clean up
    iii. trivial forwarding supports added in codegen runtime
    iv. added factory methods support in codegen

Commits that's in this PR from the devel branch:

```
7117a7e patching nvfuser conv cudnn test numerics mismatch (#2048)
65af1a4 Inserting sync for redundant parallel types is already done at the (#2023)
6ac74d1 Fix sync map (#2047)
f5bca33 Bank conflict checker improvements (#2032)
d2ca7e3 Minor update on cp.async code generation. (#1901)
d36cf61 Test file cleanup (#2040)
0b8e83f Allow non-root trivial reductions (#2037)
a2dfe40 Fix vectorize size calculation (#2035)
e040676 Use withPredicate to replace setPredicate to maintain Exprs immutable (#2025)
197221b removing ci workflow (#2034)
40e2703 Reduction rand like patch (#2031)
bc77266 Add utility for checking bank conflict of shared memory (#2029)
ddd1cf7 Add back FusionReductionWithTrivialReduction_CUDA (#2030)
fbd97e5 Revert "Cleanup trivial reduction workarounds (#2006)" (#2024)
bca20c1 Cleanup trivial reduction workarounds (#2006)
e4b6585 Trivial forwarding (#1995)
1a0e355 Fix contiguity analysis of predicates to match updated contiguity. (#1991)
a4effa6 Enable output allocation cache (#2010)
35440b7 Patching bn inference (#2016)
0f9f0b4 Add matmul benchmark (#2007)
45045cd Enable tests previously disabled due to an aliasing bug (#2005)
967aa77 Contiguous indexing for View operations (#1990)
a43cb20 Make inlining even more modular (#2004)
dc45835 Test util cleanup (#2003)
3ca21eb More strict validation (#2000)
a7a7d57 Fix build problem (#1999)
fc235b0 Just fixes comments (#1998)
482386c cleanup (#1997)
4cbe0db Improve divisible split detection (#1970)
42ccc52 Minor build fix. (#1996)
fcf8c09 Cleanup of lower_utils.cpp: Isolate out GpuLower usage (#1989)
15f2f6d Move ConcretizedBroadcastDomains to shared_ptr in GpuLower. (#1988)
8f1c7f5 Minor cleanup lower_unroll.cpp (#1994)
1d9858c Minor cleanup (#1992)
f262d9c Add support for uniform RNG (#1986)
eb1dad1 Remove non-const functions, remove GpuLower instance on build, pass in ca_map. (#1987)
634820c Add support for some empty fusion (#1981)
eabe8d8 Segment self mapping fusions (#1954)
e96aacf Enable Transpose operation (#1882)
425dce2 Add a null scheduler that helps segmenting away no-op schedules (#1835)
306d4a6 Fix canScheduleCompileTime check of transpose scheduler (#1969)
b1bd32c Minor fix (#1967)
bd93578 Enable transpose scheduler (#1927)
b7a206e Move scheduler vectorize utilities into their own file (#1959)
d9420e4 View scheduling (#1928)
c668e13 Upstream push ci fixes (#1965)
c40202b Fix dump effective bandwidth (#1962)
93505bc WAR on index mapping when exact and permissive maps differ (#1960)
45e95fd Allow splitting inner-most ID to create virtual innermost ID in transpose scheduler (#1930)
a3ecb33 Improve the comments at the beginning of index_compute.h (#1946)
f7bc341 Remove unused variables (#1955)
df3393a Some cleanup (#1957)
7d1d7c8 TVDomainGuard factory (#1953)
357ba22 Fill allocation with nan on tests (#1956)
8eafc54 Fix detection of unmappable root domains (#1952)
90a51f2 Some indexing cleanups, Add eye support (#1940)
ddc01e4 Exclude unsupported data types (#1951)
992e17c test the groups the same order as they are merged (#1949)
208262b Move detection of self mapping IDs to IterDomainGraph from (#1941)
ac4de38 Merge pull request #1945 from csarofeen/master_merge_0828
6310948 Add full, full_like, zeros, zeros_like, ones, ones_like (#1943)
aab10bc Merge remote-tracking branch 'upstream/viable/strict' into HEAD
4c254c0 Fix arange when step is negative (#1942)
89330aa Tensor factories must set the output shape as its input (#1939)
```

RUN_TORCHBENCH: nvfuser

Differential Revision: [D40869846](https://our.internmc.facebook.com/intern/diff/D40869846)
Pull Request resolved: pytorch#87779
Approved by: https://github.com/davidberard98

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment