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.

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

Unroll warp-specialized loops #3547

Merged
merged 1 commit into from
Dec 10, 2024
Merged

Unroll warp-specialized loops #3547

merged 1 commit into from
Dec 10, 2024

Conversation

zasdfgbnm
Copy link
Collaborator

@zasdfgbnm zasdfgbnm commented Dec 9, 2024

When used with #3545, this contribute a speedup of 5% of cuBLAS!

Perf together with #3545 on H100:

 Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)                                                  Name

 --------  ---------------  ---------  --------  --------  --------  --------  -----------  ----------------------------------------------------------------------------------------------------
     33.8           136319          1  136319.0  136319.0    136319    136319          0.0  <unnamed>::nvfuser_none_f0_c0_r0_g0(<unnamed>::Tensor<<unnamed>::__half, (int)3, (int)3>, <unnamed>…
     22.7            91487          1   91487.0   91487.0     91487     91487          0.0  nvjet_hsh_128x256_64x4_2x1_v_bz_coopA_NTN

nvFuser/cuBLAS: 67%

Note that the above test is run with smem epilogue disabled. I will run a test with everything combined later. Also note that this number is on H100, which is different from the H200 in #3279.

@zasdfgbnm zasdfgbnm marked this pull request as ready for review December 9, 2024 23:10
@zasdfgbnm
Copy link
Collaborator Author

!test

@zasdfgbnm zasdfgbnm requested review from naoyam and rdspring1 December 9, 2024 23:11
Copy link
Collaborator

@naoyam naoyam left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants